From 575d6ace8d5ab7c0227305f0e9f6ee0df69ee5b2 Mon Sep 17 00:00:00 2001
From: KepingYan <keping.yan@intel.com>
Date: Thu, 10 Jul 2025 07:15:15 +0000
Subject: [PATCH 1/3] [SYCL] Port to event free APIs for kernel submission

---
 src/comm/SYCLHelpers.h | 28 ++++++++++++++++------------
 1 file changed, 16 insertions(+), 12 deletions(-)

diff --git a/src/comm/SYCLHelpers.h b/src/comm/SYCLHelpers.h
index 48390229f4..954c8bb378 100644
--- a/src/comm/SYCLHelpers.h
+++ b/src/comm/SYCLHelpers.h
@@ -50,8 +50,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 +82,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 +98,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 +115,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 +134,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) \

From e1f24574d87c8dc69883fae2607416741c28fb92 Mon Sep 17 00:00:00 2001
From: KepingYan <keping.yan@intel.com>
Date: Thu, 17 Jul 2025 07:35:39 +0000
Subject: [PATCH 2/3] add build assertion

---
 src/comm/SYCLHelpers.h | 6 ++++++
 1 file changed, 6 insertions(+)

diff --git a/src/comm/SYCLHelpers.h b/src/comm/SYCLHelpers.h
index 954c8bb378..83b19c6690 100644
--- a/src/comm/SYCLHelpers.h
+++ b/src/comm/SYCLHelpers.h
@@ -3,6 +3,12 @@
 #include <comm/Scalar.h>
 #include <sycl/sycl.hpp>
 
+#if !defined(SYCL_EXT_ONEAPI_ENQUEUE_FUNCTIONS) || \
+    (SYCL_EXT_ONEAPI_ENQUEUE_FUNCTIONS != 1)
+#error This change enables event-less APIs, \
+  please ensure SYCL_EXT_ONEAPI_ENQUEUE_FUNCTIONS=1.
+#endif
+
 // sycl access address space
 static constexpr auto sycl_priv_space =
     sycl::access::address_space::private_space;

From 3d498b6c5fff300ae7dfd6d76cfde8dde540979c Mon Sep 17 00:00:00 2001
From: KepingYan <keping.yan@intel.com>
Date: Fri, 18 Jul 2025 02:13:12 +0000
Subject: [PATCH 3/3] address comment

---
 src/comm/SYCLHelpers.h | 7 +++----
 1 file changed, 3 insertions(+), 4 deletions(-)

diff --git a/src/comm/SYCLHelpers.h b/src/comm/SYCLHelpers.h
index 83b19c6690..ac67f0a99a 100644
--- a/src/comm/SYCLHelpers.h
+++ b/src/comm/SYCLHelpers.h
@@ -3,10 +3,9 @@
 #include <comm/Scalar.h>
 #include <sycl/sycl.hpp>
 
-#if !defined(SYCL_EXT_ONEAPI_ENQUEUE_FUNCTIONS) || \
-    (SYCL_EXT_ONEAPI_ENQUEUE_FUNCTIONS != 1)
-#error This change enables event-less APIs, \
-  please ensure SYCL_EXT_ONEAPI_ENQUEUE_FUNCTIONS=1.
+#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