diff --git a/src/comm/SYCLHelpers.h b/src/comm/SYCLHelpers.h
index 48390229f4..ac67f0a99a 100644
--- a/src/comm/SYCLHelpers.h
+++ b/src/comm/SYCLHelpers.h
@@ -3,6 +3,11 @@
 #include <comm/Scalar.h>
 #include <sycl/sycl.hpp>
 
+#ifndef SYCL_EXT_ONEAPI_ENQUEUE_FUNCTIONS
+#error FATAL ERROR: The SYCL_EXT_ONEAPI_ENQUEUE_FUNCTIONS is excepted to be defined \
+  to enqueue kernel without returning event. But the undef is detected!
+#endif
+
 // sycl access address space
 static constexpr auto sycl_priv_space =
     sycl::access::address_space::private_space;
@@ -50,8 +55,10 @@ static inline void sycl_kernel_submit(
     ::sycl::range<dim> range,
     ::sycl::queue q,
     ker_t ker) {
-  auto cgf = [&](::sycl::handler& cgh) { cgh.parallel_for<ker_t>(range, ker); };
-  q.submit(cgf);
+  auto cgf = [&](::sycl::handler& cgh) {
+    ::sycl::ext::oneapi::experimental::parallel_for<ker_t>(cgh, range, ker);
+  };
+  ::sycl::ext::oneapi::experimental::submit(q, cgf);
 }
 
 // Additional convention of SYCL kernel configuration. Besides construct kernel
@@ -80,10 +87,10 @@ sycl_kernel_submit(
     ker_t ker) {
   auto cgf = [&](::sycl::handler& cgh) {
     ker.sycl_ker_config_convention(cgh);
-    cgh.parallel_for<ker_t>(
-        ::sycl::nd_range<dim>(global_range, local_range), ker);
+    ::sycl::ext::oneapi::experimental::nd_launch<ker_t>(
+        cgh, ::sycl::nd_range<dim>(global_range, local_range), ker);
   };
-  q.submit(cgf);
+  ::sycl::ext::oneapi::experimental::submit(q, cgf);
 }
 
 template <typename ker_t, int dim>
@@ -96,10 +103,10 @@ sycl_kernel_submit(
     ::sycl::queue q,
     ker_t ker) {
   auto cgf = [&](::sycl::handler& cgh) {
-    cgh.parallel_for<ker_t>(
-        ::sycl::nd_range<dim>(global_range, local_range), ker);
+    ::sycl::ext::oneapi::experimental::nd_launch<ker_t>(
+        cgh, ::sycl::nd_range<dim>(global_range, local_range), ker);
   };
-  q.submit(cgf);
+  ::sycl::ext::oneapi::experimental::submit(q, cgf);
 }
 
 template <typename ker_t>
@@ -113,12 +120,13 @@ sycl_kernel_submit(
     ker_t ker) {
   auto cgf = [&](::sycl::handler& cgh) {
     ker.sycl_ker_config_convention(cgh);
-    cgh.parallel_for<ker_t>(
+    ::sycl::ext::oneapi::experimental::nd_launch<ker_t>(
+        cgh,
         ::sycl::nd_range<1>(
             ::sycl::range<1>(global_range), ::sycl::range<1>(local_range)),
         ker);
   };
-  q.submit(cgf);
+  ::sycl::ext::oneapi::experimental::submit(q, cgf);
 }
 
 template <typename ker_t>
@@ -131,12 +139,13 @@ sycl_kernel_submit(
     ::sycl::queue q,
     ker_t ker) {
   auto cgf = [&](::sycl::handler& cgh) {
-    cgh.parallel_for<ker_t>(
+    ::sycl::ext::oneapi::experimental::nd_launch<ker_t>(
+        cgh,
         ::sycl::nd_range<1>(
             ::sycl::range<1>(global_range), ::sycl::range<1>(local_range)),
         ker);
   };
-  q.submit(cgf);
+  ::sycl::ext::oneapi::experimental::submit(q, cgf);
 }
 
 #define SYCL_KERNEL_STRING(var, str) \