diff --git a/sycl/include/sycl/ext/oneapi/experimental/enqueue_functions.hpp b/sycl/include/sycl/ext/oneapi/experimental/enqueue_functions.hpp
index 8c8488a99e354..e8d1b2be337a0 100644
--- a/sycl/include/sycl/ext/oneapi/experimental/enqueue_functions.hpp
+++ b/sycl/include/sycl/ext/oneapi/experimental/enqueue_functions.hpp
@@ -263,8 +263,15 @@ template <typename KernelName = sycl::detail::auto_name, int Dimensions,
           typename KernelType, typename... ReductionsT>
 void nd_launch(handler &CGH, nd_range<Dimensions> Range,
                const KernelType &KernelObj, ReductionsT &&...Reductions) {
-  CGH.parallel_for<KernelName>(Range, std::forward<ReductionsT>(Reductions)...,
-                               KernelObj);
+  if constexpr (ext::oneapi::experimental::detail::HasKernelPropertiesGetMethod<
+                    const KernelType &>::value) {
+    CGH.parallel_for<KernelName>(
+        Range, KernelObj.get(ext::oneapi::experimental::properties_tag{}),
+        std::forward<ReductionsT>(Reductions)..., KernelObj);
+  } else {
+    CGH.parallel_for<KernelName>(
+        Range, std::forward<ReductionsT>(Reductions)..., KernelObj);
+  }
 }
 
 template <typename KernelName = sycl::detail::auto_name, int Dimensions,
diff --git a/sycl/test/check_device_code/extensions/properties/properties_kernel_device_has.cpp b/sycl/test/check_device_code/extensions/properties/properties_kernel_device_has.cpp
index 3d1c528744afd..018115a0ef2b5 100644
--- a/sycl/test/check_device_code/extensions/properties/properties_kernel_device_has.cpp
+++ b/sycl/test/check_device_code/extensions/properties/properties_kernel_device_has.cpp
@@ -27,6 +27,37 @@ static constexpr auto device_has_all = device_has<
     aspect::usm_host_allocations, aspect::usm_shared_allocations,
     aspect::ext_intel_free_memory, aspect::ext_intel_device_id>;
 
+struct TestKernelHasDevice {
+  void operator()() const {}
+  auto get(properties_tag) const { return properties{device_has_all}; }
+};
+
+struct TestKernelHasDevice_id1 {
+  void operator()(id<1>) const {}
+  auto get(properties_tag) const { return properties{device_has_all}; }
+};
+
+struct TestKernelHasDevice_id1_1 {
+  template <typename T1> void operator()(id<1>, T1 &) const {}
+  auto get(properties_tag) const { return properties{device_has_all}; }
+};
+
+struct TestKernelHasDevice_nd_item1 {
+  void operator()(nd_item<1>) const {}
+  auto get(properties_tag) const { return properties{device_has_all}; }
+};
+
+struct TestKernelHasDevice_nd_item1_1 {
+  template <typename T1> void operator()(nd_item<1>, T1 &) const {}
+  auto get(properties_tag) const { return properties{device_has_all}; }
+};
+
+struct TestKernelHasDevice_nd_item1_2 {
+  template <typename T1, typename T2>
+  void operator()(nd_item<1>, T1 &, T2 &) const {}
+  auto get(properties_tag) const { return properties{device_has_all}; }
+};
+
 int main() {
   queue Q;
   event Ev;
@@ -40,18 +71,18 @@ int main() {
   auto Redu2 = reduction<float>(nullptr, multiplies<float>());
 
   // CHECK-IR: spir_kernel void @{{.*}}WGSizeKernel0(){{.*}} #[[DHAttr1:[0-9]+]]
-  Q.single_task<class WGSizeKernel0>(Props, []() {});
+  Q.single_task<class WGSizeKernel0>(TestKernelHasDevice{});
   // CHECK-IR: spir_kernel void @{{.*}}WGSizeKernel1(){{.*}} #[[DHAttr1]]
-  Q.single_task<class WGSizeKernel1>(Ev, Props, []() {});
+  Q.single_task<class WGSizeKernel1>(Ev, TestKernelHasDevice{});
   // CHECK-IR: spir_kernel void @{{.*}}WGSizeKernel2(){{.*}} #[[DHAttr1]]
-  Q.single_task<class WGSizeKernel2>({Ev}, Props, []() {});
+  Q.single_task<class WGSizeKernel2>({Ev}, TestKernelHasDevice{});
 
   // CHECK-IR: spir_kernel void @{{.*}}WGSizeKernel3(){{.*}} #[[DHAttr2:[0-9]+]]
-  Q.parallel_for<class WGSizeKernel3>(R1, Props, [](id<1>) {});
+  Q.parallel_for<class WGSizeKernel3>(R1, TestKernelHasDevice_id1{});
   // CHECK-IR: spir_kernel void @{{.*}}WGSizeKernel4(){{.*}} #[[DHAttr2]]
-  Q.parallel_for<class WGSizeKernel4>(R1, Ev, Props, [](id<1>) {});
+  Q.parallel_for<class WGSizeKernel4>(R1, Ev, TestKernelHasDevice_id1{});
   // CHECK-IR: spir_kernel void @{{.*}}WGSizeKernel5(){{.*}} #[[DHAttr2]]
-  Q.parallel_for<class WGSizeKernel5>(R1, {Ev}, Props, [](id<1>) {});
+  Q.parallel_for<class WGSizeKernel5>(R1, {Ev}, TestKernelHasDevice_id1{});
 
   // CHECK-IR: spir_kernel void @{{.*}}MainKrn{{.*}}WGSizeKernel6{{.*}}{{.*}} #[[DHAttr2:[0-9]+]]
   Q.parallel_for<class WGSizeKernel6>(R1, Props, Redu1, [](id<1>, auto &) {});
@@ -70,57 +101,62 @@ int main() {
   Q.parallel_for<class WGSizeKernel11>(NDR1, {Ev}, Props, [](nd_item<1>) {});
 
   // CHECK-IR: spir_kernel void @{{.*}}MainKrn{{.*}}WGSizeKernel12{{.*}}{{.*}} #[[DHAttr2]]
-  Q.parallel_for<class WGSizeKernel12>(NDR1, Props, Redu1,
-                                       [](nd_item<1>, auto &) {});
+  nd_launch<class WGSizeKernel12>(Q, NDR1, TestKernelHasDevice_nd_item1_1{},
+                                  Redu1);
   // CHECK-IR: spir_kernel void @{{.*}}MainKrn{{.*}}WGSizeKernel13{{.*}}{{.*}} #[[DHAttr2]]
-  Q.parallel_for<class WGSizeKernel13>(NDR1, Ev, Props, Redu1,
-                                       [](nd_item<1>, auto &) {});
+  Q.submit([&](sycl::handler &CGH) {
+    CGH.depends_on(Ev);
+    nd_launch<class WGSizeKernel13>(CGH, NDR1, TestKernelHasDevice_nd_item1_1{},
+                                    Redu1);
+  });
   // CHECK-IR: spir_kernel void @{{.*}}MainKrn{{.*}}WGSizeKernel14{{.*}}{{.*}} #[[DHAttr2]]
-  Q.parallel_for<class WGSizeKernel14>(NDR1, {Ev}, Props, Redu1,
-                                       [](nd_item<1>, auto &) {});
+  Q.submit([&](sycl::handler &CGH) {
+    CGH.depends_on({Ev});
+    nd_launch<class WGSizeKernel14>(CGH, NDR1, TestKernelHasDevice_nd_item1_1{},
+                                    Redu1);
+  });
 
   // CHECK-IR: spir_kernel void @{{.*}}MainKrn{{.*}}WGSizeKernel15{{.*}}{{.*}} #[[DHAttr2]]
-  Q.parallel_for<class WGSizeKernel15>(NDR1, Props, Redu1, Redu2,
-                                       [](nd_item<1>, auto &, auto &) {});
+  nd_launch<class WGSizeKernel15>(Q, NDR1, TestKernelHasDevice_nd_item1_2{},
+                                  Redu1, Redu2);
   // CHECK-IR: spir_kernel void @{{.*}}MainKrn{{.*}}WGSizeKernel16{{.*}}{{.*}} #[[DHAttr2]]
-  Q.parallel_for<class WGSizeKernel16>(NDR1, Ev, Props, Redu1, Redu2,
-                                       [](nd_item<1>, auto &, auto &) {});
+  Q.submit([&](sycl::handler &CGH) {
+    CGH.depends_on(Ev);
+    nd_launch<class WGSizeKernel16>(CGH, NDR1, TestKernelHasDevice_nd_item1_2{},
+                                    Redu1, Redu2);
+  });
   // CHECK-IR: spir_kernel void @{{.*}}MainKrn{{.*}}WGSizeKernel17{{.*}}{{.*}} #[[DHAttr2]]
-  Q.parallel_for<class WGSizeKernel17>(NDR1, {Ev}, Props, Redu1, Redu2,
-                                       [](nd_item<1>, auto &, auto &) {});
+  Q.submit([&](sycl::handler &CGH) {
+    CGH.depends_on({Ev});
+    nd_launch<class WGSizeKernel17>(CGH, NDR1, TestKernelHasDevice_nd_item1_2{},
+                                    Redu1, Redu2);
+  });
 
   // CHECK-IR: spir_kernel void @{{.*}}WGSizeKernel18(){{.*}} #[[DHAttr1]]
   Q.submit([&](handler &CGH) {
-    CGH.single_task<class WGSizeKernel18>(Props, []() {});
+    CGH.single_task<class WGSizeKernel18>(TestKernelHasDevice{});
   });
 
   // CHECK-IR: spir_kernel void @{{.*}}WGSizeKernel19(){{.*}} #[[DHAttr2]]
   Q.submit([&](handler &CGH) {
-    CGH.parallel_for<class WGSizeKernel19>(R1, Props, [](id<1>) {});
+    CGH.parallel_for<class WGSizeKernel19>(R1, TestKernelHasDevice_id1{});
   });
 
   // CHECK-IR: spir_kernel void @{{.*}}MainKrn{{.*}}WGSizeKernel20{{.*}}{{.*}} #[[DHAttr2]]
   Q.submit([&](handler &CGH) {
-    CGH.parallel_for<class WGSizeKernel20>(R1, Props, Redu1,
-                                           [](id<1>, auto &) {});
+    CGH.parallel_for<class WGSizeKernel20>(R1, Props, Redu1, [](id<1>, auto &) {
+    }); // note: this one still doesn't work
   });
 
   // CHECK-IR: spir_kernel void @{{.*}}WGSizeKernel21(){{.*}} #[[DHAttr2]]
   Q.submit([&](handler &CGH) {
-    CGH.parallel_for<class WGSizeKernel21>(NDR1, Props, [](nd_item<1>) {});
+    CGH.parallel_for<class WGSizeKernel21>(NDR1,
+                                           TestKernelHasDevice_nd_item1{});
   });
 
-  // CHECK-IR: spir_kernel void @{{.*}}MainKrn{{.*}}WGSizeKernel22{{.*}}{{.*}} #[[DHAttr2]]
-  Q.submit([&](handler &CGH) {
-    CGH.parallel_for<class WGSizeKernel22>(NDR1, Props, Redu1,
-                                           [](nd_item<1>, auto &) {});
-  });
+  // DUPLICATE, REMOVED
 
-  // CHECK-IR: spir_kernel void @{{.*}}MainKrn{{.*}}WGSizeKernel23{{.*}}{{.*}} #[[DHAttr2]]
-  Q.submit([&](handler &CGH) {
-    CGH.parallel_for<class WGSizeKernel23>(NDR1, Props, Redu1, Redu2,
-                                           [](nd_item<1>, auto &, auto &) {});
-  });
+  // DUPLICATE, REMOVED
 
   // CHECK-IR: spir_kernel void @{{.*}}WGSizeKernel24(){{.*}} #[[DHAttr2]]
   Q.submit([&](handler &CGH) {