From 0598f755d0ad46778cea176c0d15f4eb68107382 Mon Sep 17 00:00:00 2001
From: "Larsen, Steffen" <steffen.larsen@intel.com>
Date: Mon, 22 Jul 2024 06:06:44 -0700
Subject: [PATCH 1/3] [SYCL][Docs] Add kernel enqueue functions for kernel and
 properties

The current version of the (sycl_ext_oneapi_kernel_properties)[https://github.com/intel/llvm/blob/sycl/sycl/doc/extensions/experimental/sycl_ext_oneapi_kernel_properties.asciidoc]
extension does not have overloads for `single_task` and `parallel_for`
that takes a `sycl::kernel` argument. This was likely omitted as all
the properties added by the extension had direct effects on how the
compiler would handle the kernel code, which cannot be done when the
kernel argument is a `sycl::kernel` object. However, as more kernel
property extension are written, some of them affect how the runtime
handles the launch of kernels and could as such be applied to such
kernel enqueue functions.

This commit adds member function overloads to `sycl::handler` for
`single_task` and `parallel_for` that takes a `sycl::kernel` argument in
the (sycl_ext_oneapi_kernel_properties)[https://github.com/intel/llvm/blob/sycl/sycl/doc/extensions/experimental/sycl_ext_oneapi_kernel_properties.asciidoc]
extension. Likewise, these new overloads are added to the implementation
of `sycl::handler`.

Signed-off-by: Larsen, Steffen <steffen.larsen@intel.com>
---
 ...sycl_ext_oneapi_kernel_properties.asciidoc |  57 +++++-
 ..._fpga_kernel_interface_properties.asciidoc |  23 ++-
 .../include/sycl/detail/kernel_properties.hpp |   1 +
 .../experimental/fpga_kernel_properties.hpp   |   9 +
 .../oneapi/kernel_properties/properties.hpp   |  13 ++
 .../sycl/ext/oneapi/properties/property.hpp   |   5 +
 sycl/include/sycl/handler.hpp                 | 187 ++++++++++++++----
 .../properties/kernel_properties_negative.cpp |  99 ++++++++++
 .../properties/properties_kernel.cpp          |   9 +
 .../properties/properties_kernel_fpga.cpp     |  11 ++
 sycl/unittests/Extensions/CMakeLists.txt      |   1 +
 .../unittests/Extensions/KernelProperties.cpp | 140 +++++++++++++
 12 files changed, 499 insertions(+), 56 deletions(-)
 create mode 100644 sycl/test/extensions/properties/kernel_properties_negative.cpp
 create mode 100644 sycl/unittests/Extensions/KernelProperties.cpp

diff --git a/sycl/doc/extensions/experimental/sycl_ext_oneapi_kernel_properties.asciidoc b/sycl/doc/extensions/experimental/sycl_ext_oneapi_kernel_properties.asciidoc
index 878f0862ac990..c03b86cc18eb7 100644
--- a/sycl/doc/extensions/experimental/sycl_ext_oneapi_kernel_properties.asciidoc
+++ b/sycl/doc/extensions/experimental/sycl_ext_oneapi_kernel_properties.asciidoc
@@ -114,6 +114,27 @@ supports.
  feature-test macro always has this value.
 |===
 
+=== Kernel Property Trait
+
+Some kernel properties carry compile-time information to affect how the compiler
+handles the corresponding kernel object. However, some of the kernel enqueue
+functions take a `sycl::kernel` object and as such it may not be possible to
+infer which kernel function it refers to at compile-time. To help identify the
+properties that have compile-time effects on the kernel, the
+`has_compile_time_kernel_effect` trait.
+
+```c++
+namespace sycl::ext::oneapi::experimental {
+
+template<typename> struct has_compile_time_kernel_effect : std::false_type {};
+
+template<typename PropertyT>
+static constexpr bool has_compile_time_kernel_effect_v =
+    has_compile_time_kernel_effect<PropertyT>::value;
+
+} // namespace sycl::ext::oneapi::experimental
+```
+
 === Kernel Properties
 
 The kernel properties below correspond to kernel attributes defined in
@@ -121,10 +142,7 @@ Section 5.8.1 of the SYCL 2020 specification.  Note that deprecated attributes
 (such as `vec_type_hint`) are not included.
 
 ```c++
-namespace sycl {
-namespace ext {
-namespace oneapi {
-namespace experimental {
+namespace sycl::ext::oneapi::experimental {
 
 // Corresponds to reqd_work_group_size
 struct work_group_size_key {
@@ -185,10 +203,16 @@ 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 experimental
-} // namespace oneapi
-} // namespace ext
-} // namespace sycl
+template <size_t... Dims>
+struct has_compile_time_kernel_effect<work_group_size_key::value_t<Dims...>> : std::true_type {};
+template <size_t... Dims>
+struct has_compile_time_kernel_effect<work_group_size_hint_key::value_t<Dims...>> : std::true_type {};
+template <uint32_t Size>
+struct has_compile_time_kernel_effect<sub_group_size_key::value_t<Size>> : std::true_type {};
+template <sycl::aspect... Aspects>
+struct has_compile_time_kernel_effect<device_has_key::value_t<Aspects...>> : std::true_type {};
+
+} // namespace sycl::ext::oneapi::experimental
 ```
 
 |===
@@ -320,6 +344,23 @@ class handler {
                                range<dimensions> workGroupSize,
                                PropertyList properties,
                                const WorkgroupFunctionType &kernelFunc);
+
+  // Available only if `!has_compile_time_kernel_effect<PropertyT>` for all
+  // properties `PropertyT` in `PropertyList`.
+  template <typename PropertyList>
+  void single_task(PropertyList properties, const kernel& kernelObject);
+
+  // Available only if `!has_compile_time_kernel_effect<PropertyT>` for all
+  // properties `PropertyT` in `PropertyList`.
+  template <int Dimensions, typename PropertyList>
+  void parallel_for(range<Dimensions> numWorkItems, PropertyList properties,
+                    const kernel& kernelObject);
+
+  // Available only if `!has_compile_time_kernel_effect<PropertyT>` for all
+  // properties `PropertyT` in `PropertyList`.
+  template <int Dimensions, typename PropertyList>
+  void parallel_for(nd_range<Dimensions> ndRange, PropertyList properties,
+                    const kernel& kernelObject);
 }
 }
 ```
diff --git a/sycl/doc/extensions/proposed/sycl_ext_intel_fpga_kernel_interface_properties.asciidoc b/sycl/doc/extensions/proposed/sycl_ext_intel_fpga_kernel_interface_properties.asciidoc
index 73bf4a6dfdb22..ae86d9e0e962b 100644
--- a/sycl/doc/extensions/proposed/sycl_ext_intel_fpga_kernel_interface_properties.asciidoc
+++ b/sycl/doc/extensions/proposed/sycl_ext_intel_fpga_kernel_interface_properties.asciidoc
@@ -92,7 +92,8 @@ using the mechanism defined in sycl_ext_oneapi_kernel_properties.
 === Kernel Interface Properties
 
 ```c++
-namespace sycl::ext::intel::experimental {
+namespace sycl::ext {
+namespace intel::experimental {
 
 enum class streaming_interface_options_enum {
   accept_downstream_stall,
@@ -167,11 +168,23 @@ inline constexpr fpga_cluster_key::value_t<
 inline constexpr fpga_cluster_key::value_t<
   fpga_cluster_options_enum::stall_free_clusters> stall_free_clusters;
 
-template <> struct is_property_key<streaming_interface_key> : std::true_type {};
-template <> struct is_property_key<register_map_interface_key> : std::true_type {};
-template <> struct is_property_key<fpga_cluster_key> : std::true_type {};
+} // intel::experimental
 
-} // namespace sycl::ext::intel::experimental
+namespace 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::fpga_cluster_key> : std::true_type {};
+
+template <intel::experimental::streaming_interface_options_enum option>
+struct has_compile_time_kernel_effect<intel::experimental::streaming_interface_key::value_t<option>> : std::true_type {};
+template <intel::experimental::register_map_interface_options_enum option>
+struct has_compile_time_kernel_effect<intel::experimental::register_map_interface_key::value_t<option>> : std::true_type {};
+template<intel::experimental::fpga_cluster_options_enum option>
+struct has_compile_time_kernel_effect<intel::experimental::fpga_cluster_key::value_t<option>> : std::true_type {};
+
+} // intel::experimental
+} // namespace sycl::ext
 ```
 
 |===
diff --git a/sycl/include/sycl/detail/kernel_properties.hpp b/sycl/include/sycl/detail/kernel_properties.hpp
index cd1b9ef7929a2..4258b594bc69a 100644
--- a/sycl/include/sycl/detail/kernel_properties.hpp
+++ b/sycl/include/sycl/detail/kernel_properties.hpp
@@ -45,6 +45,7 @@ struct PropertyMetaInfo<sycl::detail::register_alloc_mode_key::value_t<Mode>> {
   static constexpr const char *name = "sycl-register-alloc-mode";
   static constexpr sycl::detail::register_alloc_mode_enum value = Mode;
 };
+
 } // namespace detail
 } // namespace ext::oneapi::experimental
 } // namespace _V1
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 7628fb40bf4ca..df29da6da4b17 100644
--- a/sycl/include/sycl/ext/intel/experimental/fpga_kernel_properties.hpp
+++ b/sycl/include/sycl/ext/intel/experimental/fpga_kernel_properties.hpp
@@ -103,6 +103,15 @@ struct is_property_key_of<
     intel::experimental::fpga_kernel_attribute<T, PropertyListT>>
     : std::true_type {};
 
+template <intel::experimental::streaming_interface_options_enum option>
+struct has_compile_time_kernel_effect<
+    intel::experimental::streaming_interface_key::value_t<option>>
+    : std::true_type {};
+template <intel::experimental::register_map_interface_options_enum option>
+struct has_compile_time_kernel_effect<
+    intel::experimental::register_map_interface_key::value_t<option>>
+    : std::true_type {};
+
 namespace detail {
 template <intel::experimental::streaming_interface_options_enum Stall_Free>
 struct PropertyMetaInfo<
diff --git a/sycl/include/sycl/ext/oneapi/kernel_properties/properties.hpp b/sycl/include/sycl/ext/oneapi/kernel_properties/properties.hpp
index e46ab88c43172..81106954df912 100644
--- a/sycl/include/sycl/ext/oneapi/kernel_properties/properties.hpp
+++ b/sycl/include/sycl/ext/oneapi/kernel_properties/properties.hpp
@@ -238,6 +238,19 @@ template <> struct is_property_key<work_group_progress_key> : std::true_type {};
 template <> struct is_property_key<sub_group_progress_key> : std::true_type {};
 template <> struct is_property_key<work_item_progress_key> : std::true_type {};
 
+template <size_t... Dims>
+struct has_compile_time_kernel_effect<work_group_size_key::value_t<Dims...>>
+    : std::true_type {};
+template <size_t... Dims>
+struct has_compile_time_kernel_effect<
+    work_group_size_hint_key::value_t<Dims...>> : std::true_type {};
+template <uint32_t Size>
+struct has_compile_time_kernel_effect<sub_group_size_key::value_t<Size>>
+    : std::true_type {};
+template <sycl::aspect... Aspects>
+struct has_compile_time_kernel_effect<device_has_key::value_t<Aspects...>>
+    : std::true_type {};
+
 namespace detail {
 
 template <size_t Dim0, size_t... Dims>
diff --git a/sycl/include/sycl/ext/oneapi/properties/property.hpp b/sycl/include/sycl/ext/oneapi/properties/property.hpp
index f8c365240a706..cedf6b1ac03dd 100644
--- a/sycl/include/sycl/ext/oneapi/properties/property.hpp
+++ b/sycl/include/sycl/ext/oneapi/properties/property.hpp
@@ -274,6 +274,11 @@ struct is_property_key
 };
 template <typename, typename> struct is_property_key_of : std::false_type {};
 
+template <typename> struct has_compile_time_kernel_effect : std::false_type {};
+template <typename PropertyT>
+static constexpr bool has_compile_time_kernel_effect_v =
+    has_compile_time_kernel_effect<PropertyT>::value;
+
 } // namespace experimental
 } // namespace oneapi
 } // namespace ext
diff --git a/sycl/include/sycl/handler.hpp b/sycl/include/sycl/handler.hpp
index 117217eb685d0..bffdccc413bf5 100644
--- a/sycl/include/sycl/handler.hpp
+++ b/sycl/include/sycl/handler.hpp
@@ -265,6 +265,18 @@ struct GetMergedKernelProperties<
       PropertiesT, get_method_properties>;
 };
 
+// Checks that none of the properties in the property list has compile-time
+// effects on the kernel.
+template <typename T>
+struct NoPropertyHasCompileTimeKernelEffect : std::false_type {};
+template <typename... Ts>
+struct NoPropertyHasCompileTimeKernelEffect<
+    ext::oneapi::experimental::detail::properties_t<Ts...>> {
+  static constexpr bool value =
+      !(ext::oneapi::experimental::has_compile_time_kernel_effect_v<Ts> ||
+        ... || false);
+};
+
 #if __SYCL_ID_QUERIES_FIT_IN_INT__
 template <typename T, typename ValT>
 typename std::enable_if_t<std::is_same<ValT, size_t>::value ||
@@ -970,28 +982,11 @@ class __SYCL_EXPORT handler {
     }
   }
 
-  /// Process kernel properties.
+  /// Process runtime kernel properties.
   ///
   /// Stores information about kernel properties into the handler.
-  template <
-      typename KernelName,
-      typename PropertiesT = ext::oneapi::experimental::empty_properties_t>
-  void processProperties(PropertiesT Props) {
-    using KI = detail::KernelInfo<KernelName>;
-    static_assert(
-        ext::oneapi::experimental::is_property_list<PropertiesT>::value,
-        "Template type is not a property list.");
-    static_assert(
-        !PropertiesT::template has_property<
-            sycl::ext::intel::experimental::fp_control_key>() ||
-            (PropertiesT::template has_property<
-                 sycl::ext::intel::experimental::fp_control_key>() &&
-             KI::isESIMD()),
-        "Floating point control property is supported for ESIMD kernels only.");
-    static_assert(
-        !PropertiesT::template has_property<
-            sycl::ext::oneapi::experimental::indirectly_callable_key>(),
-        "indirectly_callable property cannot be applied to SYCL kernels");
+  template <typename PropertiesT>
+  void processLaunchProperties(PropertiesT Props) {
     if constexpr (PropertiesT::template has_property<
                       sycl::ext::intel::experimental::cache_config_key>()) {
       auto Config = Props.template get_property<
@@ -1042,6 +1037,32 @@ class __SYCL_EXPORT handler {
     checkAndSetClusterRange(Props);
   }
 
+  /// Process kernel properties.
+  ///
+  /// Stores information about kernel properties into the handler.
+  template <
+      typename KernelName,
+      typename PropertiesT = ext::oneapi::experimental::empty_properties_t>
+  void processProperties(PropertiesT Props) {
+    using KI = detail::KernelInfo<KernelName>;
+    static_assert(
+        ext::oneapi::experimental::is_property_list<PropertiesT>::value,
+        "Template type is not a property list.");
+    static_assert(
+        !PropertiesT::template has_property<
+            sycl::ext::intel::experimental::fp_control_key>() ||
+            (PropertiesT::template has_property<
+                 sycl::ext::intel::experimental::fp_control_key>() &&
+             KI::isESIMD()),
+        "Floating point control property is supported for ESIMD kernels only.");
+    static_assert(
+        !PropertiesT::template has_property<
+            sycl::ext::oneapi::experimental::indirectly_callable_key>(),
+        "indirectly_callable property cannot be applied to SYCL kernels");
+
+    processLaunchProperties(Props);
+  }
+
   /// Checks whether it is possible to copy the source shape to the destination
   /// shape(the shapes are described by the accessor ranges) by using
   /// copying by regions of memory and not copying element by element
@@ -1440,12 +1461,15 @@ class __SYCL_EXPORT handler {
   ///
   /// \param NumWorkItems is a range defining indexing space.
   /// \param Kernel is a SYCL kernel function.
-  template <int Dims>
-  void parallel_for_impl(range<Dims> NumWorkItems, kernel Kernel) {
+  /// \param Properties is the properties.
+  template <int Dims, typename PropertiesT>
+  void parallel_for_impl(range<Dims> NumWorkItems, PropertiesT Props,
+                         kernel Kernel) {
     throwIfActionIsCreated();
     MKernel = detail::getSyclObjImpl(std::move(Kernel));
     detail::checkValueRange<Dims>(NumWorkItems);
     setNDRangeDescriptor(std::move(NumWorkItems));
+    processLaunchProperties<PropertiesT>(Props);
     setType(detail::CGType::Kernel);
     setNDRangeUsed(false);
     extractArgsAndReqs();
@@ -2125,28 +2149,22 @@ class __SYCL_EXPORT handler {
   ///
   /// \param Kernel is a SYCL kernel object.
   void single_task(kernel Kernel) {
-    throwIfActionIsCreated();
-    // Ignore any set kernel bundles and use the one associated with the kernel
-    setHandlerKernelBundle(Kernel);
-    // No need to check if range is out of INT_MAX limits as it's compile-time
-    // known constant
-    setNDRangeDescriptor(range<1>{1});
-    MKernel = detail::getSyclObjImpl(std::move(Kernel));
-    setType(detail::CGType::Kernel);
-    extractArgsAndReqs();
-    MKernelName = getKernelName();
+    single_task(ext::oneapi::experimental::empty_properties_t{}, Kernel);
   }
 
   void parallel_for(range<1> NumWorkItems, kernel Kernel) {
-    parallel_for_impl(NumWorkItems, Kernel);
+    parallel_for_impl(NumWorkItems,
+                      ext::oneapi::experimental::empty_properties_t{}, Kernel);
   }
 
   void parallel_for(range<2> NumWorkItems, kernel Kernel) {
-    parallel_for_impl(NumWorkItems, Kernel);
+    parallel_for_impl(NumWorkItems,
+                      ext::oneapi::experimental::empty_properties_t{}, Kernel);
   }
 
   void parallel_for(range<3> NumWorkItems, kernel Kernel) {
-    parallel_for_impl(NumWorkItems, Kernel);
+    parallel_for_impl(NumWorkItems,
+                      ext::oneapi::experimental::empty_properties_t{}, Kernel);
   }
 
   /// Defines and invokes a SYCL kernel function for the specified range and
@@ -2180,14 +2198,8 @@ class __SYCL_EXPORT handler {
   /// well as offset.
   /// \param Kernel is a SYCL kernel function.
   template <int Dims> void parallel_for(nd_range<Dims> NDRange, kernel Kernel) {
-    throwIfActionIsCreated();
-    MKernel = detail::getSyclObjImpl(std::move(Kernel));
-    detail::checkValueRange<Dims>(NDRange);
-    setNDRangeDescriptor(std::move(NDRange));
-    setType(detail::CGType::Kernel);
-    setNDRangeUsed(true);
-    extractArgsAndReqs();
-    MKernelName = getKernelName();
+    parallel_for(NDRange, ext::oneapi::experimental::empty_properties_t{},
+                 Kernel);
   }
 
   /// Defines and invokes a SYCL kernel function.
@@ -2573,6 +2585,95 @@ class __SYCL_EXPORT handler {
         NumWorkGroups, WorkGroupSize, Props, KernelFunc);
   }
 
+  /// Invokes a SYCL kernel.
+  ///
+  /// Executes exactly once. The kernel invocation method has no functors and
+  /// cannot be called on host.
+  ///
+  /// \param Kernel is a SYCL kernel object.
+  /// \param Props is the properties for the launch.
+  template <typename PropertiesT>
+  std::enable_if_t<ext::oneapi::experimental::is_property_list_v<PropertiesT>,
+                   void>
+  single_task(PropertiesT Props, kernel Kernel) {
+    static_assert(
+        detail::NoPropertyHasCompileTimeKernelEffect<PropertiesT>::value,
+        "This kernel enqueue function does not allow properties with "
+        "compile-time kernel effects.");
+    throwIfActionIsCreated();
+    // Ignore any set kernel bundles and use the one associated with the kernel
+    setHandlerKernelBundle(Kernel);
+    // No need to check if range is out of INT_MAX limits as it's compile-time
+    // known constant
+    setNDRangeDescriptor(range<1>{1});
+    processLaunchProperties(Props);
+    MKernel = detail::getSyclObjImpl(std::move(Kernel));
+    setType(detail::CGType::Kernel);
+    extractArgsAndReqs();
+    MKernelName = getKernelName();
+  }
+
+  template <typename PropertiesT>
+  std::enable_if_t<ext::oneapi::experimental::is_property_list_v<PropertiesT>,
+                   void>
+  parallel_for(range<1> NumWorkItems, PropertiesT Props, kernel Kernel) {
+    static_assert(
+        detail::NoPropertyHasCompileTimeKernelEffect<PropertiesT>::value,
+        "This kernel enqueue function does not allow properties with "
+        "compile-time kernel effects.");
+    parallel_for_impl(NumWorkItems, Props, Kernel);
+  }
+
+  template <typename PropertiesT>
+  std::enable_if_t<ext::oneapi::experimental::is_property_list_v<PropertiesT>,
+                   void>
+  parallel_for(range<2> NumWorkItems, PropertiesT Props, kernel Kernel) {
+    static_assert(
+        detail::NoPropertyHasCompileTimeKernelEffect<PropertiesT>::value,
+        "This kernel enqueue function does not allow properties with "
+        "compile-time kernel effects.");
+    parallel_for_impl(NumWorkItems, Props, Kernel);
+  }
+
+  template <typename PropertiesT>
+  std::enable_if_t<ext::oneapi::experimental::is_property_list_v<PropertiesT>,
+                   void>
+  parallel_for(range<3> NumWorkItems, PropertiesT Props, kernel Kernel) {
+    static_assert(
+        detail::NoPropertyHasCompileTimeKernelEffect<PropertiesT>::value,
+        "This kernel enqueue function does not allow properties with "
+        "compile-time kernel effects.");
+    parallel_for_impl(NumWorkItems, Props, Kernel);
+  }
+
+  /// Defines and invokes a SYCL kernel function for the specified range and
+  /// offsets.
+  ///
+  /// The SYCL kernel function is defined as SYCL kernel object.
+  ///
+  /// \param NDRange is a ND-range defining global and local sizes as
+  /// well as offset.
+  /// \param Props is the properties for the launch.
+  /// \param Kernel is a SYCL kernel function.
+  template <int Dims, typename PropertiesT>
+  std::enable_if_t<ext::oneapi::experimental::is_property_list_v<PropertiesT>,
+                   void>
+  parallel_for(nd_range<Dims> NDRange, PropertiesT Props, kernel Kernel) {
+    static_assert(
+        detail::NoPropertyHasCompileTimeKernelEffect<PropertiesT>::value,
+        "This kernel enqueue function does not allow properties with "
+        "compile-time kernel effects.");
+    throwIfActionIsCreated();
+    MKernel = detail::getSyclObjImpl(std::move(Kernel));
+    detail::checkValueRange<Dims>(NDRange);
+    setNDRangeDescriptor(std::move(NDRange));
+    processLaunchProperties(Props);
+    setType(detail::CGType::Kernel);
+    setNDRangeUsed(true);
+    extractArgsAndReqs();
+    MKernelName = getKernelName();
+  }
+
   // Clean up KERNELFUNC macro.
 #undef _KERNELFUNCPARAM
 
diff --git a/sycl/test/extensions/properties/kernel_properties_negative.cpp b/sycl/test/extensions/properties/kernel_properties_negative.cpp
new file mode 100644
index 0000000000000..965b8e62e4c36
--- /dev/null
+++ b/sycl/test/extensions/properties/kernel_properties_negative.cpp
@@ -0,0 +1,99 @@
+// RUN: %clangxx -ferror-limit=0 %fsycl-host-only -fsyntax-only -Xclang -verify -Xclang -verify-ignore-unexpected=note,warning %s
+
+// Negative tests for kernel properties.
+
+#include <sycl/sycl.hpp>
+
+namespace oneapi = sycl::ext::oneapi::experimental;
+
+extern sycl::kernel TestKernel;
+
+int main() {
+  sycl::queue Q{};
+
+  oneapi::properties props1{oneapi::sub_group_size<8>};
+  oneapi::properties props2{
+      oneapi::sub_group_size<8>,
+      oneapi::work_group_progress<oneapi::forward_progress_guarantee::parallel,
+                                  oneapi::execution_scope::root_group>};
+
+  Q.submit([&](sycl::handler &CGH) {
+    // expected-error-re@sycl/handler.hpp:* {{static assertion failed due to requirement {{.*}} This kernel enqueue function does not allow properties with compile-time kernel effects.}}
+    CGH.single_task(props1, TestKernel);
+  });
+
+  Q.submit([&](sycl::handler &CGH) {
+    // expected-error-re@sycl/handler.hpp:* {{static assertion failed due to requirement {{.*}} This kernel enqueue function does not allow properties with compile-time kernel effects.}}
+    CGH.single_task(props2, TestKernel);
+  });
+
+  Q.submit([&](sycl::handler &CGH) {
+    // expected-error-re@sycl/handler.hpp:* {{static assertion failed due to requirement {{.*}} This kernel enqueue function does not allow properties with compile-time kernel effects.}}
+    CGH.parallel_for(sycl::range<1>{1}, props1, TestKernel);
+  });
+
+  Q.submit([&](sycl::handler &CGH) {
+    // expected-error-re@sycl/handler.hpp:* {{static assertion failed due to requirement {{.*}} This kernel enqueue function does not allow properties with compile-time kernel effects.}}
+    CGH.parallel_for(sycl::range<1>{1}, props2, TestKernel);
+  });
+
+  Q.submit([&](sycl::handler &CGH) {
+    // expected-error-re@sycl/handler.hpp:* {{static assertion failed due to requirement {{.*}} This kernel enqueue function does not allow properties with compile-time kernel effects.}}
+    CGH.parallel_for(sycl::range<2>{1, 1}, props1, TestKernel);
+  });
+
+  Q.submit([&](sycl::handler &CGH) {
+    // expected-error-re@sycl/handler.hpp:* {{static assertion failed due to requirement {{.*}} This kernel enqueue function does not allow properties with compile-time kernel effects.}}
+    CGH.parallel_for(sycl::range<2>{1, 1}, props2, TestKernel);
+  });
+
+  Q.submit([&](sycl::handler &CGH) {
+    // expected-error-re@sycl/handler.hpp:* {{static assertion failed due to requirement {{.*}} This kernel enqueue function does not allow properties with compile-time kernel effects.}}
+    CGH.parallel_for(sycl::range<3>{1, 1, 1}, props1, TestKernel);
+  });
+
+  Q.submit([&](sycl::handler &CGH) {
+    // expected-error-re@sycl/handler.hpp:* {{static assertion failed due to requirement {{.*}} This kernel enqueue function does not allow properties with compile-time kernel effects.}}
+    CGH.parallel_for(sycl::range<3>{1, 1, 1}, props2, TestKernel);
+  });
+
+  Q.submit([&](sycl::handler &CGH) {
+    // expected-error-re@sycl/handler.hpp:* {{static assertion failed due to requirement {{.*}} This kernel enqueue function does not allow properties with compile-time kernel effects.}}
+    CGH.parallel_for(sycl::nd_range<1>{sycl::range<1>{1}, sycl::range<1>{1}},
+                     props1, TestKernel);
+  });
+
+  Q.submit([&](sycl::handler &CGH) {
+    // expected-error-re@sycl/handler.hpp:* {{static assertion failed due to requirement {{.*}} This kernel enqueue function does not allow properties with compile-time kernel effects.}}
+    CGH.parallel_for(sycl::nd_range<1>{sycl::range<1>{1}, sycl::range<1>{1}},
+                     props2, TestKernel);
+  });
+
+  Q.submit([&](sycl::handler &CGH) {
+    // expected-error-re@sycl/handler.hpp:* {{static assertion failed due to requirement {{.*}} This kernel enqueue function does not allow properties with compile-time kernel effects.}}
+    CGH.parallel_for(
+        sycl::nd_range<2>{sycl::range<2>{1, 1}, sycl::range<2>{1, 1}}, props1,
+        TestKernel);
+  });
+
+  Q.submit([&](sycl::handler &CGH) {
+    // expected-error-re@sycl/handler.hpp:* {{static assertion failed due to requirement {{.*}} This kernel enqueue function does not allow properties with compile-time kernel effects.}}
+    CGH.parallel_for(
+        sycl::nd_range<2>{sycl::range<2>{1, 1}, sycl::range<2>{1, 1}}, props2,
+        TestKernel);
+  });
+
+  Q.submit([&](sycl::handler &CGH) {
+    // expected-error-re@sycl/handler.hpp:* {{static assertion failed due to requirement {{.*}} This kernel enqueue function does not allow properties with compile-time kernel effects.}}
+    CGH.parallel_for(
+        sycl::nd_range<3>{sycl::range<3>{1, 1, 1}, sycl::range<3>{1, 1, 1}},
+        props1, TestKernel);
+  });
+
+  Q.submit([&](sycl::handler &CGH) {
+    // expected-error-re@sycl/handler.hpp:* {{static assertion failed due to requirement {{.*}} This kernel enqueue function does not allow properties with compile-time kernel effects.}}
+    CGH.parallel_for(
+        sycl::nd_range<3>{sycl::range<3>{1, 1, 1}, sycl::range<3>{1, 1, 1}},
+        props2, TestKernel);
+  });
+}
diff --git a/sycl/test/extensions/properties/properties_kernel.cpp b/sycl/test/extensions/properties/properties_kernel.cpp
index 3868c23f7535c..14a4ca1cee330 100644
--- a/sycl/test/extensions/properties/properties_kernel.cpp
+++ b/sycl/test/extensions/properties/properties_kernel.cpp
@@ -42,6 +42,15 @@ int main() {
   static_assert(is_property_key<sub_group_size_key>::value);
   static_assert(is_property_key<device_has_key>::value);
 
+  static_assert(
+      has_compile_time_kernel_effect<work_group_size_key::value_t<1>>::value);
+  static_assert(has_compile_time_kernel_effect<
+                work_group_size_hint_key::value_t<1>>::value);
+  static_assert(
+      has_compile_time_kernel_effect<sub_group_size_key::value_t<28>>::value);
+  static_assert(has_compile_time_kernel_effect<
+                device_has_key::value_t<aspect::fp64>>::value);
+
   static_assert(is_property_value<decltype(work_group_size<1>)>::value);
   static_assert(is_property_value<decltype(work_group_size<2, 2>)>::value);
   static_assert(is_property_value<decltype(work_group_size<3, 3, 3>)>::value);
diff --git a/sycl/test/extensions/properties/properties_kernel_fpga.cpp b/sycl/test/extensions/properties/properties_kernel_fpga.cpp
index 2acd767c386ae..e578492bc1ae3 100644
--- a/sycl/test/extensions/properties/properties_kernel_fpga.cpp
+++ b/sycl/test/extensions/properties/properties_kernel_fpga.cpp
@@ -16,6 +16,17 @@ int main() {
   static_assert(oneapi::experimental::is_property_key<
                 intel::experimental::pipelined_key>::value);
 
+  // Check that oneapi::experimental::has_compile_time_kernel_effect is
+  // correctly specialized
+  static_assert(oneapi::experimental::has_compile_time_kernel_effect<
+                intel::experimental::register_map_interface_key::value_t<
+                    intel::experimental::register_map_interface_options_enum::
+                        wait_for_done_write>>::value);
+  static_assert(oneapi::experimental::has_compile_time_kernel_effect<
+                intel::experimental::register_map_interface_key::value_t<
+                    intel::experimental::register_map_interface_options_enum::
+                        wait_for_done_write>>::value);
+
   // Check that oneapi::experimental::is_property_value is correctly specialized
   static_assert(oneapi::experimental::is_property_value<
                 decltype(intel::experimental::streaming_interface<
diff --git a/sycl/unittests/Extensions/CMakeLists.txt b/sycl/unittests/Extensions/CMakeLists.txt
index 285e2a5ebd675..b276d509fcc6a 100644
--- a/sycl/unittests/Extensions/CMakeLists.txt
+++ b/sycl/unittests/Extensions/CMakeLists.txt
@@ -13,6 +13,7 @@ add_sycl_unittest(ExtensionsTests OBJECT
   EnqueueFunctionsEvents.cpp
   DiscardEvent.cpp
   ProfilingTag.cpp
+  KernelProperties.cpp
 )
 
 add_subdirectory(CommandGraph)
diff --git a/sycl/unittests/Extensions/KernelProperties.cpp b/sycl/unittests/Extensions/KernelProperties.cpp
new file mode 100644
index 0000000000000..ffd38ee6bed13
--- /dev/null
+++ b/sycl/unittests/Extensions/KernelProperties.cpp
@@ -0,0 +1,140 @@
+//==----------------------- KernelProperties.cpp ---------------------------==//
+//
+// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
+// See https://llvm.org/LICENSE.txt for license information.
+// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
+//
+//===----------------------------------------------------------------------===//
+
+#include <helpers/PiMock.hpp>
+#include <helpers/TestKernel.hpp>
+
+#include <gtest/gtest.h>
+
+#include <sycl/detail/core.hpp>
+#include <sycl/kernel_bundle.hpp>
+
+namespace oneapiext = sycl::ext::oneapi::experimental;
+
+namespace {
+
+inline pi_result after_piKernelGetInfo(pi_kernel kernel,
+                                       pi_kernel_info param_name,
+                                       size_t param_value_size,
+                                       void *param_value,
+                                       size_t *param_value_size_ret) {
+  constexpr char MockKernel[] = "TestKernel";
+  if (param_name == PI_KERNEL_INFO_FUNCTION_NAME) {
+    if (param_value) {
+      assert(param_value_size == sizeof(MockKernel));
+      std::memcpy(param_value, MockKernel, sizeof(MockKernel));
+    }
+    if (param_value_size_ret)
+      *param_value_size_ret = sizeof(MockKernel);
+  }
+  return PI_SUCCESS;
+}
+
+class KernelPropertiesTests : public ::testing::Test {
+public:
+  KernelPropertiesTests()
+      : Mock{}, Q{sycl::context(Mock.getPlatform()), sycl::default_selector_v} {
+  }
+
+  inline sycl::kernel GetTestKernel() {
+    auto KB = sycl::get_kernel_bundle<sycl::bundle_state::executable>(
+        Q.get_context());
+    return KB.get_kernel<TestKernel<>>();
+  }
+
+  template <typename FuncT> void RunForwardProgressTest(const FuncT &F) {
+    sycl::kernel K = GetTestKernel();
+
+    try {
+      F(K);
+      FAIL() << "Exception was expected.";
+    } catch (sycl::exception &E) {
+      // Intention of this test is to ensure that properties without kernel
+      // effects are applied correctly through the interface. For forward
+      // progress we can do this by expecting that the kernel reports that it is
+      // using an unsupported forward progress.
+      ASSERT_EQ(E.code(),
+                sycl::make_error_code(sycl::errc::feature_not_supported));
+    }
+  }
+
+protected:
+  void SetUp() override {
+    Mock.redefineAfter<sycl::detail::PiApiKind::piKernelGetInfo>(
+        after_piKernelGetInfo);
+  }
+
+  sycl::unittest::PiMock Mock;
+  sycl::queue Q;
+};
+
+oneapiext::properties ForwardProgressProp{oneapiext::work_group_progress<
+    oneapiext::forward_progress_guarantee::parallel,
+    oneapiext::execution_scope::root_group>};
+
+TEST_F(KernelPropertiesTests, SingleTaskKernelObjForwardProgress) {
+  RunForwardProgressTest([&](sycl::kernel &K) {
+    Q.submit(
+        [&](sycl::handler &CGH) { CGH.single_task(ForwardProgressProp, K); });
+  });
+}
+
+TEST_F(KernelPropertiesTests, ParallelForRange1DKernelObjForwardProgress) {
+  RunForwardProgressTest([&](sycl::kernel &K) {
+    Q.submit([&](sycl::handler &CGH) {
+      CGH.parallel_for(sycl::range<1>{1}, ForwardProgressProp, K);
+    });
+  });
+}
+
+TEST_F(KernelPropertiesTests, ParallelForRange2DKernelObjForwardProgress) {
+  RunForwardProgressTest([&](sycl::kernel &K) {
+    Q.submit([&](sycl::handler &CGH) {
+      CGH.parallel_for(sycl::range<2>{1, 1}, ForwardProgressProp, K);
+    });
+  });
+}
+
+TEST_F(KernelPropertiesTests, ParallelForRange3DKernelObjForwardProgress) {
+  RunForwardProgressTest([&](sycl::kernel &K) {
+    Q.submit([&](sycl::handler &CGH) {
+      CGH.parallel_for(sycl::range<3>{1, 1, 1}, ForwardProgressProp, K);
+    });
+  });
+}
+
+TEST_F(KernelPropertiesTests, ParallelForNDRange1DKernelObjForwardProgress) {
+  RunForwardProgressTest([&](sycl::kernel &K) {
+    Q.submit([&](sycl::handler &CGH) {
+      CGH.parallel_for(sycl::nd_range<1>{sycl::range<1>{1}, sycl::range<1>{1}},
+                       ForwardProgressProp, K);
+    });
+  });
+}
+
+TEST_F(KernelPropertiesTests, ParallelForNDRange2DKernelObjForwardProgress) {
+  RunForwardProgressTest([&](sycl::kernel &K) {
+    Q.submit([&](sycl::handler &CGH) {
+      CGH.parallel_for(
+          sycl::nd_range<2>{sycl::range<2>{1, 1}, sycl::range<2>{1, 1}},
+          ForwardProgressProp, K);
+    });
+  });
+}
+
+TEST_F(KernelPropertiesTests, ParallelForNDRange3DKernelObjForwardProgress) {
+  RunForwardProgressTest([&](sycl::kernel &K) {
+    Q.submit([&](sycl::handler &CGH) {
+      CGH.parallel_for(
+          sycl::nd_range<3>{sycl::range<3>{1, 1, 1}, sycl::range<3>{1, 1, 1}},
+          ForwardProgressProp, K);
+    });
+  });
+}
+
+} // namespace

From 4a7765a46f227a32d84babcb084e10a3ab03d9eb Mon Sep 17 00:00:00 2001
From: "Larsen, Steffen" <steffen.larsen@intel.com>
Date: Tue, 23 Jul 2024 07:35:41 -0700
Subject: [PATCH 2/3] Don't add new trait to spec

Signed-off-by: Larsen, Steffen <steffen.larsen@intel.com>
---
 ...sycl_ext_oneapi_kernel_properties.asciidoc | 42 +++----------------
 ..._fpga_kernel_interface_properties.asciidoc |  7 ----
 .../experimental/fpga_kernel_properties.hpp   |  6 +--
 .../oneapi/kernel_properties/properties.hpp   | 12 +++---
 .../sycl/ext/oneapi/properties/property.hpp   |  7 +---
 sycl/include/sycl/handler.hpp                 |  2 +-
 .../properties/properties_kernel.cpp          | 12 +++---
 .../properties/properties_kernel_fpga.cpp     |  6 +--
 8 files changed, 27 insertions(+), 67 deletions(-)

diff --git a/sycl/doc/extensions/experimental/sycl_ext_oneapi_kernel_properties.asciidoc b/sycl/doc/extensions/experimental/sycl_ext_oneapi_kernel_properties.asciidoc
index c03b86cc18eb7..6101bbfc2b4c7 100644
--- a/sycl/doc/extensions/experimental/sycl_ext_oneapi_kernel_properties.asciidoc
+++ b/sycl/doc/extensions/experimental/sycl_ext_oneapi_kernel_properties.asciidoc
@@ -114,27 +114,6 @@ supports.
  feature-test macro always has this value.
 |===
 
-=== Kernel Property Trait
-
-Some kernel properties carry compile-time information to affect how the compiler
-handles the corresponding kernel object. However, some of the kernel enqueue
-functions take a `sycl::kernel` object and as such it may not be possible to
-infer which kernel function it refers to at compile-time. To help identify the
-properties that have compile-time effects on the kernel, the
-`has_compile_time_kernel_effect` trait.
-
-```c++
-namespace sycl::ext::oneapi::experimental {
-
-template<typename> struct has_compile_time_kernel_effect : std::false_type {};
-
-template<typename PropertyT>
-static constexpr bool has_compile_time_kernel_effect_v =
-    has_compile_time_kernel_effect<PropertyT>::value;
-
-} // namespace sycl::ext::oneapi::experimental
-```
-
 === Kernel Properties
 
 The kernel properties below correspond to kernel attributes defined in
@@ -203,15 +182,6 @@ 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 {};
 
-template <size_t... Dims>
-struct has_compile_time_kernel_effect<work_group_size_key::value_t<Dims...>> : std::true_type {};
-template <size_t... Dims>
-struct has_compile_time_kernel_effect<work_group_size_hint_key::value_t<Dims...>> : std::true_type {};
-template <uint32_t Size>
-struct has_compile_time_kernel_effect<sub_group_size_key::value_t<Size>> : std::true_type {};
-template <sycl::aspect... Aspects>
-struct has_compile_time_kernel_effect<device_has_key::value_t<Aspects...>> : std::true_type {};
-
 } // namespace sycl::ext::oneapi::experimental
 ```
 
@@ -345,19 +315,19 @@ class handler {
                                PropertyList properties,
                                const WorkgroupFunctionType &kernelFunc);
 
-  // Available only if `!has_compile_time_kernel_effect<PropertyT>` for all
-  // properties `PropertyT` in `PropertyList`.
+  // Available only if all properties in `PropertList` only have launch-related
+  // effects.
   template <typename PropertyList>
   void single_task(PropertyList properties, const kernel& kernelObject);
 
-  // Available only if `!has_compile_time_kernel_effect<PropertyT>` for all
-  // properties `PropertyT` in `PropertyList`.
+  // Available only if all properties in `PropertList` only have launch-related
+  // effects.
   template <int Dimensions, typename PropertyList>
   void parallel_for(range<Dimensions> numWorkItems, PropertyList properties,
                     const kernel& kernelObject);
 
-  // Available only if `!has_compile_time_kernel_effect<PropertyT>` for all
-  // properties `PropertyT` in `PropertyList`.
+  // Available only if all properties in `PropertList` only have launch-related
+  // effects.
   template <int Dimensions, typename PropertyList>
   void parallel_for(nd_range<Dimensions> ndRange, PropertyList properties,
                     const kernel& kernelObject);
diff --git a/sycl/doc/extensions/proposed/sycl_ext_intel_fpga_kernel_interface_properties.asciidoc b/sycl/doc/extensions/proposed/sycl_ext_intel_fpga_kernel_interface_properties.asciidoc
index ae86d9e0e962b..13cea0f399603 100644
--- a/sycl/doc/extensions/proposed/sycl_ext_intel_fpga_kernel_interface_properties.asciidoc
+++ b/sycl/doc/extensions/proposed/sycl_ext_intel_fpga_kernel_interface_properties.asciidoc
@@ -176,13 +176,6 @@ template <> struct is_property_key<intel::experimental::streaming_interface_key>
 template <> struct is_property_key<intel::experimental::register_map_interface_key> : std::true_type {};
 template <> struct is_property_key<intel::experimental::fpga_cluster_key> : std::true_type {};
 
-template <intel::experimental::streaming_interface_options_enum option>
-struct has_compile_time_kernel_effect<intel::experimental::streaming_interface_key::value_t<option>> : std::true_type {};
-template <intel::experimental::register_map_interface_options_enum option>
-struct has_compile_time_kernel_effect<intel::experimental::register_map_interface_key::value_t<option>> : std::true_type {};
-template<intel::experimental::fpga_cluster_options_enum option>
-struct has_compile_time_kernel_effect<intel::experimental::fpga_cluster_key::value_t<option>> : std::true_type {};
-
 } // intel::experimental
 } // namespace sycl::ext
 ```
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 df29da6da4b17..85b125aa127cd 100644
--- a/sycl/include/sycl/ext/intel/experimental/fpga_kernel_properties.hpp
+++ b/sycl/include/sycl/ext/intel/experimental/fpga_kernel_properties.hpp
@@ -103,16 +103,16 @@ struct is_property_key_of<
     intel::experimental::fpga_kernel_attribute<T, PropertyListT>>
     : std::true_type {};
 
+namespace detail {
 template <intel::experimental::streaming_interface_options_enum option>
-struct has_compile_time_kernel_effect<
+struct HasCompileTimeEffect<
     intel::experimental::streaming_interface_key::value_t<option>>
     : std::true_type {};
 template <intel::experimental::register_map_interface_options_enum option>
-struct has_compile_time_kernel_effect<
+struct HasCompileTimeEffect<
     intel::experimental::register_map_interface_key::value_t<option>>
     : std::true_type {};
 
-namespace detail {
 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/oneapi/kernel_properties/properties.hpp b/sycl/include/sycl/ext/oneapi/kernel_properties/properties.hpp
index 81106954df912..47e4338db25ee 100644
--- a/sycl/include/sycl/ext/oneapi/kernel_properties/properties.hpp
+++ b/sycl/include/sycl/ext/oneapi/kernel_properties/properties.hpp
@@ -238,21 +238,21 @@ template <> struct is_property_key<work_group_progress_key> : std::true_type {};
 template <> struct is_property_key<sub_group_progress_key> : std::true_type {};
 template <> struct is_property_key<work_item_progress_key> : std::true_type {};
 
+namespace detail {
+
 template <size_t... Dims>
-struct has_compile_time_kernel_effect<work_group_size_key::value_t<Dims...>>
+struct HasCompileTimeEffect<work_group_size_key::value_t<Dims...>>
     : std::true_type {};
 template <size_t... Dims>
-struct has_compile_time_kernel_effect<
+struct HasCompileTimeEffect<
     work_group_size_hint_key::value_t<Dims...>> : std::true_type {};
 template <uint32_t Size>
-struct has_compile_time_kernel_effect<sub_group_size_key::value_t<Size>>
+struct HasCompileTimeEffect<sub_group_size_key::value_t<Size>>
     : std::true_type {};
 template <sycl::aspect... Aspects>
-struct has_compile_time_kernel_effect<device_has_key::value_t<Aspects...>>
+struct HasCompileTimeEffect<device_has_key::value_t<Aspects...>>
     : std::true_type {};
 
-namespace detail {
-
 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/properties/property.hpp b/sycl/include/sycl/ext/oneapi/properties/property.hpp
index cedf6b1ac03dd..f196242b808ba 100644
--- a/sycl/include/sycl/ext/oneapi/properties/property.hpp
+++ b/sycl/include/sycl/ext/oneapi/properties/property.hpp
@@ -266,6 +266,8 @@ template <typename PropertyT> struct PropertyMetaInfo {
   static constexpr std::nullptr_t value = nullptr;
 };
 
+template <typename> struct HasCompileTimeEffect : std::false_type {};
+
 } // namespace detail
 
 template <typename T>
@@ -274,11 +276,6 @@ struct is_property_key
 };
 template <typename, typename> struct is_property_key_of : std::false_type {};
 
-template <typename> struct has_compile_time_kernel_effect : std::false_type {};
-template <typename PropertyT>
-static constexpr bool has_compile_time_kernel_effect_v =
-    has_compile_time_kernel_effect<PropertyT>::value;
-
 } // namespace experimental
 } // namespace oneapi
 } // namespace ext
diff --git a/sycl/include/sycl/handler.hpp b/sycl/include/sycl/handler.hpp
index bffdccc413bf5..5e99c6d9c74a5 100644
--- a/sycl/include/sycl/handler.hpp
+++ b/sycl/include/sycl/handler.hpp
@@ -273,7 +273,7 @@ template <typename... Ts>
 struct NoPropertyHasCompileTimeKernelEffect<
     ext::oneapi::experimental::detail::properties_t<Ts...>> {
   static constexpr bool value =
-      !(ext::oneapi::experimental::has_compile_time_kernel_effect_v<Ts> ||
+      !(ext::oneapi::experimental::detail::HasCompileTimeEffect<Ts>::value ||
         ... || false);
 };
 
diff --git a/sycl/test/extensions/properties/properties_kernel.cpp b/sycl/test/extensions/properties/properties_kernel.cpp
index 14a4ca1cee330..dfad5f9c638e9 100644
--- a/sycl/test/extensions/properties/properties_kernel.cpp
+++ b/sycl/test/extensions/properties/properties_kernel.cpp
@@ -42,13 +42,13 @@ int main() {
   static_assert(is_property_key<sub_group_size_key>::value);
   static_assert(is_property_key<device_has_key>::value);
 
-  static_assert(
-      has_compile_time_kernel_effect<work_group_size_key::value_t<1>>::value);
-  static_assert(has_compile_time_kernel_effect<
+  static_assert(sycl::ext::oneapi::experimental::detail::HasCompileTimeEffect<
+                work_group_size_key::value_t<1>>::value);
+  static_assert(sycl::ext::oneapi::experimental::detail::HasCompileTimeEffect<
                 work_group_size_hint_key::value_t<1>>::value);
-  static_assert(
-      has_compile_time_kernel_effect<sub_group_size_key::value_t<28>>::value);
-  static_assert(has_compile_time_kernel_effect<
+  static_assert(sycl::ext::oneapi::experimental::detail::HasCompileTimeEffect<
+                sub_group_size_key::value_t<28>>::value);
+  static_assert(sycl::ext::oneapi::experimental::detail::HasCompileTimeEffect<
                 device_has_key::value_t<aspect::fp64>>::value);
 
   static_assert(is_property_value<decltype(work_group_size<1>)>::value);
diff --git a/sycl/test/extensions/properties/properties_kernel_fpga.cpp b/sycl/test/extensions/properties/properties_kernel_fpga.cpp
index e578492bc1ae3..da3d0e8aa2637 100644
--- a/sycl/test/extensions/properties/properties_kernel_fpga.cpp
+++ b/sycl/test/extensions/properties/properties_kernel_fpga.cpp
@@ -16,13 +16,13 @@ int main() {
   static_assert(oneapi::experimental::is_property_key<
                 intel::experimental::pipelined_key>::value);
 
-  // Check that oneapi::experimental::has_compile_time_kernel_effect is
+  // Check that oneapi::experimental::detail::HasCompileTimeEffect is
   // correctly specialized
-  static_assert(oneapi::experimental::has_compile_time_kernel_effect<
+  static_assert(oneapi::experimental::detail::HasCompileTimeEffect<
                 intel::experimental::register_map_interface_key::value_t<
                     intel::experimental::register_map_interface_options_enum::
                         wait_for_done_write>>::value);
-  static_assert(oneapi::experimental::has_compile_time_kernel_effect<
+  static_assert(oneapi::experimental::detail::HasCompileTimeEffect<
                 intel::experimental::register_map_interface_key::value_t<
                     intel::experimental::register_map_interface_options_enum::
                         wait_for_done_write>>::value);

From c730e9d890288801d2c9395c6b9b89b922545bb3 Mon Sep 17 00:00:00 2001
From: "Larsen, Steffen" <steffen.larsen@intel.com>
Date: Tue, 23 Jul 2024 07:46:07 -0700
Subject: [PATCH 3/3] Fix formatting

Signed-off-by: Larsen, Steffen <steffen.larsen@intel.com>
---
 sycl/include/sycl/ext/oneapi/kernel_properties/properties.hpp | 4 ++--
 1 file changed, 2 insertions(+), 2 deletions(-)

diff --git a/sycl/include/sycl/ext/oneapi/kernel_properties/properties.hpp b/sycl/include/sycl/ext/oneapi/kernel_properties/properties.hpp
index 47e4338db25ee..cd805dcd8e082 100644
--- a/sycl/include/sycl/ext/oneapi/kernel_properties/properties.hpp
+++ b/sycl/include/sycl/ext/oneapi/kernel_properties/properties.hpp
@@ -244,8 +244,8 @@ template <size_t... Dims>
 struct HasCompileTimeEffect<work_group_size_key::value_t<Dims...>>
     : std::true_type {};
 template <size_t... Dims>
-struct HasCompileTimeEffect<
-    work_group_size_hint_key::value_t<Dims...>> : std::true_type {};
+struct HasCompileTimeEffect<work_group_size_hint_key::value_t<Dims...>>
+    : std::true_type {};
 template <uint32_t Size>
 struct HasCompileTimeEffect<sub_group_size_key::value_t<Size>>
     : std::true_type {};