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..85b125aa127cd 100644
--- a/sycl/include/sycl/ext/intel/experimental/fpga_kernel_properties.hpp
+++ b/sycl/include/sycl/ext/intel/experimental/fpga_kernel_properties.hpp
@@ -104,6 +104,15 @@ struct is_property_key_of<
     : std::true_type {};
 
 namespace detail {
+template <intel::experimental::streaming_interface_options_enum option>
+struct HasCompileTimeEffect<
+    intel::experimental::streaming_interface_key::value_t<option>>
+    : std::true_type {};
+template <intel::experimental::register_map_interface_options_enum option>
+struct HasCompileTimeEffect<
+    intel::experimental::register_map_interface_key::value_t<option>>
+    : 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/oneapi/experimental/enqueue_functions.hpp b/sycl/include/sycl/ext/oneapi/experimental/enqueue_functions.hpp
index 7d463d7c66eca..ad78970d82709 100644
--- a/sycl/include/sycl/ext/oneapi/experimental/enqueue_functions.hpp
+++ b/sycl/include/sycl/ext/oneapi/experimental/enqueue_functions.hpp
@@ -34,6 +34,16 @@ template <typename RangeT>
 constexpr bool is_range_or_nd_range_v = is_range_or_nd_range<RangeT>::value;
 
 template <typename LCRangeT, typename LCPropertiesT> struct LaunchConfigAccess;
+
+// 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<properties_t<Ts...>> {
+  static constexpr bool value =
+      !(HasCompileTimeEffect<Ts>::value || ... || false);
+};
 } // namespace detail
 
 // Available only when Range is range or nd_range
@@ -42,6 +52,11 @@ template <
     typename = std::enable_if_t<
         ext::oneapi::experimental::detail::is_range_or_nd_range_v<RangeT>>>
 class launch_config {
+  static_assert(ext::oneapi::experimental::detail::
+                    NoPropertyHasCompileTimeKernelEffect<PropertiesT>::value,
+                "launch_config does not allow properties with compile-time "
+                "kernel effects.");
+
 public:
   launch_config(RangeT Range, PropertiesT Properties = {})
       : MRange{Range}, MProperties{Properties} {}
@@ -187,7 +202,8 @@ void parallel_for(handler &CGH,
                                                         Properties>
       ConfigAccess(Config);
   CGH.set_args<ArgsT...>(std::forward<ArgsT>(Args)...);
-  CGH.parallel_for(ConfigAccess.getRange(), KernelObj);
+  sycl::detail::HandlerAccess::parallelForImpl(
+      CGH, ConfigAccess.getRange(), ConfigAccess.getProperties(), KernelObj);
 }
 
 template <int Dimensions, typename Properties, typename... ArgsT>
@@ -263,7 +279,8 @@ void nd_launch(handler &CGH,
                                                         Properties>
       ConfigAccess(Config);
   CGH.set_args<ArgsT...>(std::forward<ArgsT>(Args)...);
-  CGH.parallel_for(ConfigAccess.getRange(), KernelObj);
+  sycl::detail::HandlerAccess::parallelForImpl(
+      CGH, ConfigAccess.getRange(), ConfigAccess.getProperties(), KernelObj);
 }
 
 template <int Dimensions, typename Properties, typename... ArgsT>
diff --git a/sycl/include/sycl/ext/oneapi/kernel_properties/properties.hpp b/sycl/include/sycl/ext/oneapi/kernel_properties/properties.hpp
index e46ab88c43172..cd805dcd8e082 100644
--- a/sycl/include/sycl/ext/oneapi/kernel_properties/properties.hpp
+++ b/sycl/include/sycl/ext/oneapi/kernel_properties/properties.hpp
@@ -240,6 +240,19 @@ template <> struct is_property_key<work_item_progress_key> : std::true_type {};
 
 namespace detail {
 
+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 {};
+template <uint32_t Size>
+struct HasCompileTimeEffect<sub_group_size_key::value_t<Size>>
+    : std::true_type {};
+template <sycl::aspect... Aspects>
+struct HasCompileTimeEffect<device_has_key::value_t<Aspects...>>
+    : 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/properties/property.hpp b/sycl/include/sycl/ext/oneapi/properties/property.hpp
index f8c365240a706..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>
diff --git a/sycl/include/sycl/handler.hpp b/sycl/include/sycl/handler.hpp
index 1e312f04ffc51..5408431a4c9db 100644
--- a/sycl/include/sycl/handler.hpp
+++ b/sycl/include/sycl/handler.hpp
@@ -970,28 +970,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 +1025,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,18 +1449,44 @@ 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();
     MKernelName = getKernelName();
   }
 
+  /// 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 Properties is the properties.
+  /// \param Kernel is a SYCL kernel function.
+  template <int Dims, typename PropertiesT>
+  void parallel_for_impl(nd_range<Dims> NDRange, PropertiesT Props,
+                         kernel Kernel) {
+    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();
+  }
+
   /// Hierarchical kernel invocation method of a kernel defined as a lambda
   /// encoding the body of each work-group to launch.
   ///
@@ -2163,15 +2198,18 @@ class __SYCL_EXPORT handler {
   }
 
   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
@@ -2205,14 +2243,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_impl(NDRange, ext::oneapi::experimental::empty_properties_t{},
+                      Kernel);
   }
 
   /// Defines and invokes a SYCL kernel function.
@@ -3741,6 +3773,12 @@ class HandlerAccess {
   static void internalProfilingTagImpl(handler &Handler) {
     Handler.internalProfilingTagImpl();
   }
+
+  template <typename RangeT, typename PropertiesT>
+  static void parallelForImpl(handler &Handler, RangeT Range, PropertiesT Props,
+                              kernel Kernel) {
+    Handler.parallel_for_impl(Range, Props, Kernel);
+  }
 };
 } // namespace detail
 
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..fdc910c92a3d2
--- /dev/null
+++ b/sycl/test/extensions/properties/kernel_properties_negative.cpp
@@ -0,0 +1,25 @@
+// 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>};
+
+  // expected-error-re@sycl/ext/oneapi/experimental/enqueue_functions.hpp:* {{static assertion failed due to requirement {{.*}} launch_config does not allow properties with compile-time kernel effects.}}
+  oneapi::launch_config<sycl::range<1>, decltype(props1)> LC1{{1}, props1};
+
+  // expected-error-re@sycl/ext/oneapi/experimental/enqueue_functions.hpp:* {{static assertion failed due to requirement {{.*}} launch_config does not allow properties with compile-time kernel effects.}}
+  oneapi::launch_config<sycl::range<1>, decltype(props2)> LC22{{1}, props2};
+}
diff --git a/sycl/test/extensions/properties/properties_kernel.cpp b/sycl/test/extensions/properties/properties_kernel.cpp
index 3868c23f7535c..dfad5f9c638e9 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(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(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);
   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..51f5fee5ce1a4 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::detail::HasCompileTimeEffect is
+  // correctly specialized
+  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::detail::HasCompileTimeEffect<
+                intel::experimental::streaming_interface_key::value_t<
+                    intel::experimental::streaming_interface_options_enum::
+                        accept_downstream_stall>>::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..31aa984726b0a
--- /dev/null
+++ b/sycl/unittests/Extensions/KernelProperties.cpp
@@ -0,0 +1,129 @@
+//==----------------------- 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/ext/oneapi/experimental/enqueue_functions.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, ParallelFor1DKernelObjForwardProgress) {
+  RunForwardProgressTest([&](sycl::kernel &K) {
+    oneapiext::launch_config<sycl::range<1>, decltype(ForwardProgressProp)> LC{
+        {1}, ForwardProgressProp};
+    oneapiext::parallel_for(Q, LC, K);
+  });
+}
+
+TEST_F(KernelPropertiesTests, ParallelFor2DKernelObjForwardProgress) {
+  RunForwardProgressTest([&](sycl::kernel &K) {
+    oneapiext::launch_config<sycl::range<2>, decltype(ForwardProgressProp)> LC{
+        {1, 1}, ForwardProgressProp};
+    oneapiext::parallel_for(Q, LC, K);
+  });
+}
+
+TEST_F(KernelPropertiesTests, ParallelFor3DKernelObjForwardProgress) {
+  RunForwardProgressTest([&](sycl::kernel &K) {
+    oneapiext::launch_config<sycl::range<3>, decltype(ForwardProgressProp)> LC{
+        {1, 1, 1}, ForwardProgressProp};
+    oneapiext::parallel_for(Q, LC, K);
+  });
+}
+
+TEST_F(KernelPropertiesTests, NDLaunch1DKernelObjForwardProgress) {
+  RunForwardProgressTest([&](sycl::kernel &K) {
+    oneapiext::launch_config<sycl::nd_range<1>, decltype(ForwardProgressProp)>
+        LC{{{1}, {1}}, ForwardProgressProp};
+    oneapiext::nd_launch(Q, LC, K);
+  });
+}
+
+TEST_F(KernelPropertiesTests, NDLaunch2DKernelObjForwardProgress) {
+  RunForwardProgressTest([&](sycl::kernel &K) {
+    oneapiext::launch_config<sycl::nd_range<2>, decltype(ForwardProgressProp)>
+        LC{{{1, 1}, {1, 1}}, ForwardProgressProp};
+    oneapiext::nd_launch(Q, LC, K);
+  });
+}
+
+TEST_F(KernelPropertiesTests, NDLaunch3DKernelObjForwardProgress) {
+  RunForwardProgressTest([&](sycl::kernel &K) {
+    oneapiext::launch_config<sycl::nd_range<3>, decltype(ForwardProgressProp)>
+        LC{{{1, 1, 1}, {1, 1, 1}}, ForwardProgressProp};
+    oneapiext::nd_launch(Q, LC, K);
+  });
+}
+
+} // namespace