diff --git a/sycl/include/sycl/handler.hpp b/sycl/include/sycl/handler.hpp
index a61c1dfff8e28..3a6528f5cb561 100644
--- a/sycl/include/sycl/handler.hpp
+++ b/sycl/include/sycl/handler.hpp
@@ -2437,6 +2437,21 @@ class __SYCL_EXPORT handler {
     parallel_for_impl<KernelName>(Range, Properties, std::move(KernelFunc));
   }
 
+  /// 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 <typename PropertiesT, int Dims>
+  void parallel_for(nd_range<Dims> NDRange, PropertiesT Properties,
+                    kernel Kernel) {
+    parallel_for_impl(NDRange, Properties, Kernel);
+  }
+
   /// Reductions @{
 
   template <typename KernelName = detail::auto_name, typename PropertiesT,
diff --git a/sycl/test-e2e/Properties/kernel.cpp b/sycl/test-e2e/Properties/kernel.cpp
new file mode 100644
index 0000000000000..37dcb5b59c10e
--- /dev/null
+++ b/sycl/test-e2e/Properties/kernel.cpp
@@ -0,0 +1,57 @@
+// REQUIRES: gpu, level_zero
+
+// RUN: %{build} -o %t.out
+// RUN: %{run} %t.out
+
+// Verifies that a kernel can be launched with properities.
+
+#include <numeric>
+#include <sycl/detail/core.hpp>
+#include <sycl/properties/all_properties.hpp>
+#include <sycl/usm.hpp>
+
+using namespace sycl;
+using namespace sycl::ext::intel::experimental;
+using namespace sycl::ext::oneapi::experimental;
+
+// TODO: remove SYCL_EXTERNAL once it is no longer needed.
+auto constexpr SYCLSource = R"===(
+#include <sycl/sycl.hpp>
+
+// use extern "C" to avoid name mangling
+extern "C" SYCL_EXTERNAL SYCL_EXT_ONEAPI_FUNCTION_PROPERTY((sycl::ext::oneapi::experimental::nd_range_kernel<1>))
+void ff_cp() {}
+)===";
+
+int main() {
+  namespace syclex = sycl::ext::oneapi::experimental;
+  using source_kb = sycl::kernel_bundle<sycl::bundle_state::ext_oneapi_source>;
+  using exe_kb = sycl::kernel_bundle<sycl::bundle_state::executable>;
+
+  queue Queue{};
+  sycl::context Ctx = Queue.get_context();
+
+  syclex::properties Properties{};
+
+  // Create from source.
+  source_kb kbSrc = syclex::create_kernel_bundle_from_source(
+      Ctx, syclex::source_language::sycl, SYCLSource, syclex::properties{});
+
+  // Compilation of empty prop list, no devices.
+  exe_kb kbExe1 = syclex::build(kbSrc);
+
+  // clang-format off
+
+  sycl::nd_range<1> R1{{10}, {1}};
+  // extern "C" was used, so the name "ff_cp" is not mangled and can be used directly.
+  sycl::kernel Kernel = kbExe1.ext_oneapi_get_kernel("ff_cp");
+
+  // clang-format on
+
+  Queue.submit([&](sycl::handler &Handler) {
+    Handler.parallel_for(R1, Properties, Kernel);
+  });
+  Queue.wait();
+
+  return 0;
+}
diff --git a/sycl/test-e2e/no_sycl_hpp_in_e2e_tests.cpp b/sycl/test-e2e/no_sycl_hpp_in_e2e_tests.cpp
index a7c2329b23dc7..5c4c5b0233463 100644
--- a/sycl/test-e2e/no_sycl_hpp_in_e2e_tests.cpp
+++ b/sycl/test-e2e/no_sycl_hpp_in_e2e_tests.cpp
@@ -7,7 +7,7 @@
 // CHECK-DAG: no_sycl_hpp_in_e2e_tests.cpp
 // CHECK-DAG: lit.cfg.py
 //
-// CHECK-NUM-MATCHES: 5
+// CHECK-NUM-MATCHES: 6
 //
 // This test verifies that `<sycl/sycl.hpp>` isn't used in E2E tests. Instead,
 // fine-grained includes should used, see