From 496b6f0f6f2ebff819b0a362a54f70db5b40c967 Mon Sep 17 00:00:00 2001
From: Hugh Bird <hugh.bird@codeplay.com>
Date: Thu, 20 Jun 2024 19:11:35 +0100
Subject: [PATCH 01/10] ROCm implementation

---
 src/dft/backends/rocfft/backward.cpp | 17 +++++++++--------
 src/dft/backends/rocfft/forward.cpp  | 22 ++++++++++++++--------
 2 files changed, 23 insertions(+), 16 deletions(-)

diff --git a/src/dft/backends/rocfft/backward.cpp b/src/dft/backends/rocfft/backward.cpp
index 5ff0e2a1f..25329bf26 100644
--- a/src/dft/backends/rocfft/backward.cpp
+++ b/src/dft/backends/rocfft/backward.cpp
@@ -17,6 +17,7 @@
 * SPDX-License-Identifier: Apache-2.0
 *******************************************************************************/
 
+#include <sycl/ext/adaptivecpp/custom_operation.hpp>
 #if __has_include(<sycl/sycl.hpp>)
 #include <sycl/sycl.hpp>
 #else
@@ -78,7 +79,7 @@ ONEMKL_EXPORT void compute_backward(descriptor_type &desc,
         auto inout_acc = inout.template get_access<sycl::access::mode::read_write>(cgh);
         commit->add_buffer_workspace_dependency_if_rqd("compute_backward", cgh);
 
-        cgh.host_task([=](sycl::interop_handle ih) {
+        cgh.AdaptiveCpp_enqueue_custom_operation([=](sycl::interop_handle ih) {
             auto stream = detail::setup_stream(func_name, ih, info);
 
             auto inout_native = reinterpret_cast<void *>(
@@ -113,7 +114,7 @@ ONEMKL_EXPORT void compute_backward(descriptor_type &desc,
         auto inout_im_acc = inout_im.template get_access<sycl::access::mode::read_write>(cgh);
         commit->add_buffer_workspace_dependency_if_rqd("compute_backward", cgh);
 
-        cgh.host_task([=](sycl::interop_handle ih) {
+        cgh.AdaptiveCpp_enqueue_custom_operation([=](sycl::interop_handle ih) {
             auto stream = detail::setup_stream(func_name, ih, info);
 
             std::array<void *, 2> inout_native{
@@ -148,7 +149,7 @@ ONEMKL_EXPORT void compute_backward(descriptor_type &desc,
         auto out_acc = out.template get_access<sycl::access::mode::read_write>(cgh);
         commit->add_buffer_workspace_dependency_if_rqd("compute_backward", cgh);
 
-        cgh.host_task([=](sycl::interop_handle ih) {
+        cgh.AdaptiveCpp_enqueue_custom_operation([=](sycl::interop_handle ih) {
             const std::string func_name = "compute_backward(desc, in, out)";
             auto stream = detail::setup_stream(func_name, ih, info);
 
@@ -184,7 +185,7 @@ ONEMKL_EXPORT void compute_backward(descriptor_type &desc,
         auto out_im_acc = out_im.template get_access<sycl::access::mode::read_write>(cgh);
         commit->add_buffer_workspace_dependency_if_rqd("compute_backward", cgh);
 
-        cgh.host_task([=](sycl::interop_handle ih) {
+        cgh.AdaptiveCpp_enqueue_custom_operation([=](sycl::interop_handle ih) {
             const std::string func_name = "compute_backward(desc, in_re, in_im, out_re, out_im)";
             auto stream = detail::setup_stream(func_name, ih, info);
 
@@ -239,7 +240,7 @@ ONEMKL_EXPORT sycl::event compute_backward(descriptor_type &desc, fwd<descriptor
         cgh.depends_on(deps);
         commit->depend_on_last_usm_workspace_event_if_rqd(cgh);
 
-        cgh.host_task([=](sycl::interop_handle ih) {
+        cgh.AdaptiveCpp_enqueue_custom_operation([=](sycl::interop_handle ih) {
             auto stream = detail::setup_stream(func_name, ih, info);
 
             void *inout_ptr = inout;
@@ -273,7 +274,7 @@ ONEMKL_EXPORT sycl::event compute_backward(descriptor_type &desc, scalar<descrip
         cgh.depends_on(deps);
         commit->depend_on_last_usm_workspace_event_if_rqd(cgh);
 
-        cgh.host_task([=](sycl::interop_handle ih) {
+        cgh.AdaptiveCpp_enqueue_custom_operation([=](sycl::interop_handle ih) {
             auto stream = detail::setup_stream(func_name, ih, info);
 
             std::array<void *, 2> inout_native{ inout_re + offsets[0], inout_im + offsets[0] };
@@ -305,7 +306,7 @@ ONEMKL_EXPORT sycl::event compute_backward(descriptor_type &desc, bwd<descriptor
         cgh.depends_on(deps);
         commit->depend_on_last_usm_workspace_event_if_rqd(cgh);
 
-        cgh.host_task([=](sycl::interop_handle ih) {
+        cgh.AdaptiveCpp_enqueue_custom_operation([=](sycl::interop_handle ih) {
             const std::string func_name = "compute_backward(desc, in, out, deps)";
             auto stream = detail::setup_stream(func_name, ih, info);
 
@@ -336,7 +337,7 @@ ONEMKL_EXPORT sycl::event compute_backward(descriptor_type &desc, scalar<descrip
         cgh.depends_on(deps);
         commit->depend_on_last_usm_workspace_event_if_rqd(cgh);
 
-        cgh.host_task([=](sycl::interop_handle ih) {
+        cgh.AdaptiveCpp_enqueue_custom_operation([=](sycl::interop_handle ih) {
             const std::string func_name =
                 "compute_backward(desc, in_re, in_im, out_re, out_im, deps)";
             auto stream = detail::setup_stream(func_name, ih, info);
diff --git a/src/dft/backends/rocfft/forward.cpp b/src/dft/backends/rocfft/forward.cpp
index 70d3d0f97..fc3fe5900 100644
--- a/src/dft/backends/rocfft/forward.cpp
+++ b/src/dft/backends/rocfft/forward.cpp
@@ -18,11 +18,17 @@
 *******************************************************************************/
 
 #include <type_traits>
+
+#include <sycl/ext/adaptivecpp/custom_operation.hpp>
 #if __has_include(<sycl/sycl.hpp>)
 #include <sycl/sycl.hpp>
 #else
 #include <CL/sycl.hpp>
 #endif
+#ifndef SYCL_EXT_ACPP_ENQUEUE_CUSTOM_OPERATION
+# error Extension not implemented!
+#endif
+
 
 #include "oneapi/mkl/exceptions.hpp"
 
@@ -81,7 +87,7 @@ ONEMKL_EXPORT void compute_forward(descriptor_type &desc,
         auto inout_acc = inout.template get_access<sycl::access::mode::read_write>(cgh);
         commit->add_buffer_workspace_dependency_if_rqd("compute_forward", cgh);
 
-        cgh.host_task([=](sycl::interop_handle ih) {
+        cgh.AdaptiveCpp_enqueue_custom_operation([=](sycl::interop_handle ih) {
             auto stream = detail::setup_stream(func_name, ih, info);
 
             auto inout_native = reinterpret_cast<void *>(
@@ -116,7 +122,7 @@ ONEMKL_EXPORT void compute_forward(descriptor_type &desc,
         auto inout_im_acc = inout_im.template get_access<sycl::access::mode::read_write>(cgh);
         commit->add_buffer_workspace_dependency_if_rqd("compute_forward", cgh);
 
-        cgh.host_task([=](sycl::interop_handle ih) {
+        cgh.AdaptiveCpp_enqueue_custom_operation([=](sycl::interop_handle ih) {
             auto stream = detail::setup_stream(func_name, ih, info);
 
             std::array<void *, 2> inout_native{
@@ -150,7 +156,7 @@ ONEMKL_EXPORT void compute_forward(descriptor_type &desc, sycl::buffer<fwd<descr
         auto out_acc = out.template get_access<sycl::access::mode::read_write>(cgh);
         commit->add_buffer_workspace_dependency_if_rqd("compute_forward", cgh);
 
-        cgh.host_task([=](sycl::interop_handle ih) {
+        cgh.AdaptiveCpp_enqueue_custom_operation([=](sycl::interop_handle ih) {
             const std::string func_name = "compute_forward(desc, in, out)";
             auto stream = detail::setup_stream(func_name, ih, info);
 
@@ -186,7 +192,7 @@ ONEMKL_EXPORT void compute_forward(descriptor_type &desc,
         auto out_im_acc = out_im.template get_access<sycl::access::mode::read_write>(cgh);
         commit->add_buffer_workspace_dependency_if_rqd("compute_forward", cgh);
 
-        cgh.host_task([=](sycl::interop_handle ih) {
+        cgh.AdaptiveCpp_enqueue_custom_operation([=](sycl::interop_handle ih) {
             const std::string func_name = "compute_forward(desc, in_re, in_im, out_re, out_im)";
             auto stream = detail::setup_stream(func_name, ih, info);
 
@@ -241,7 +247,7 @@ ONEMKL_EXPORT sycl::event compute_forward(descriptor_type &desc, fwd<descriptor_
         cgh.depends_on(deps);
         commit->depend_on_last_usm_workspace_event_if_rqd(cgh);
 
-        cgh.host_task([=](sycl::interop_handle ih) {
+        cgh.AdaptiveCpp_enqueue_custom_operation([=](sycl::interop_handle ih) {
             auto stream = detail::setup_stream(func_name, ih, info);
 
             void *inout_ptr = inout;
@@ -274,7 +280,7 @@ ONEMKL_EXPORT sycl::event compute_forward(descriptor_type &desc, scalar<descript
     sycl::event sycl_event = queue.submit([&](sycl::handler &cgh) {
         cgh.depends_on(deps);
         commit->depend_on_last_usm_workspace_event_if_rqd(cgh);
-        cgh.host_task([=](sycl::interop_handle ih) {
+        cgh.AdaptiveCpp_enqueue_custom_operation([=](sycl::interop_handle ih) {
             auto stream = detail::setup_stream(func_name, ih, info);
 
             std::array<void *, 2> inout_native{ inout_re + offsets[0], inout_im + offsets[0] };
@@ -306,7 +312,7 @@ ONEMKL_EXPORT sycl::event compute_forward(descriptor_type &desc, fwd<descriptor_
         cgh.depends_on(deps);
         commit->depend_on_last_usm_workspace_event_if_rqd(cgh);
 
-        cgh.host_task([=](sycl::interop_handle ih) {
+        cgh.AdaptiveCpp_enqueue_custom_operation([=](sycl::interop_handle ih) {
             const std::string func_name = "compute_forward(desc, in, out, deps)";
             auto stream = detail::setup_stream(func_name, ih, info);
 
@@ -337,7 +343,7 @@ ONEMKL_EXPORT sycl::event compute_forward(descriptor_type &desc, scalar<descript
         cgh.depends_on(deps);
         commit->depend_on_last_usm_workspace_event_if_rqd(cgh);
 
-        cgh.host_task([=](sycl::interop_handle ih) {
+        cgh.AdaptiveCpp_enqueue_custom_operation([=](sycl::interop_handle ih) {
             const std::string func_name =
                 "compute_forward(desc, in_re, in_im, out_re, out_im, deps)";
             auto stream = detail::setup_stream(func_name, ih, info);

From 8c5479b886738a78434ca724edcd1b2a1efb7d84 Mon Sep 17 00:00:00 2001
From: Hugh Bird <hugh.bird@codeplay.com>
Date: Wed, 26 Jun 2024 22:11:01 +0100
Subject: [PATCH 02/10] Remove host synchronization

---
 src/dft/backends/rocfft/backward.cpp       | 16 ++++++++--------
 src/dft/backends/rocfft/execute_helper.hpp | 20 ++++++++------------
 src/dft/backends/rocfft/forward.cpp        | 16 ++++++++--------
 3 files changed, 24 insertions(+), 28 deletions(-)

diff --git a/src/dft/backends/rocfft/backward.cpp b/src/dft/backends/rocfft/backward.cpp
index 25329bf26..46b44d724 100644
--- a/src/dft/backends/rocfft/backward.cpp
+++ b/src/dft/backends/rocfft/backward.cpp
@@ -86,7 +86,7 @@ ONEMKL_EXPORT void compute_backward(descriptor_type &desc,
                 reinterpret_cast<fwd<descriptor_type> *>(detail::native_mem(ih, inout_acc)) +
                 offsets[0]);
             detail::execute_checked(func_name, plan, &inout_native, nullptr, info);
-            detail::sync_checked(func_name, stream);
+            
         });
     });
 }
@@ -126,7 +126,7 @@ ONEMKL_EXPORT void compute_backward(descriptor_type &desc,
                                          offsets[0])
             };
             detail::execute_checked(func_name, plan, inout_native.data(), nullptr, info);
-            detail::sync_checked(func_name, stream);
+            
         });
     });
 }
@@ -160,7 +160,7 @@ ONEMKL_EXPORT void compute_backward(descriptor_type &desc,
                 reinterpret_cast<fwd<descriptor_type> *>(detail::native_mem(ih, out_acc)) +
                 offsets[1]);
             detail::execute_checked(func_name, plan, &in_native, &out_native, info);
-            detail::sync_checked(func_name, stream);
+            
         });
     });
 }
@@ -206,7 +206,7 @@ ONEMKL_EXPORT void compute_backward(descriptor_type &desc,
                                          offsets[1])
             };
             detail::execute_checked(func_name, plan, in_native.data(), out_native.data(), info);
-            detail::sync_checked(func_name, stream);
+            
         });
     });
 }
@@ -245,7 +245,7 @@ ONEMKL_EXPORT sycl::event compute_backward(descriptor_type &desc, fwd<descriptor
 
             void *inout_ptr = inout;
             detail::execute_checked(func_name, plan, &inout_ptr, nullptr, info);
-            detail::sync_checked(func_name, stream);
+            
         });
     });
     commit->set_last_usm_workspace_event_if_rqd(sycl_event);
@@ -279,7 +279,7 @@ ONEMKL_EXPORT sycl::event compute_backward(descriptor_type &desc, scalar<descrip
 
             std::array<void *, 2> inout_native{ inout_re + offsets[0], inout_im + offsets[0] };
             detail::execute_checked(func_name, plan, inout_native.data(), nullptr, info);
-            detail::sync_checked(func_name, stream);
+            
         });
     });
     commit->set_last_usm_workspace_event_if_rqd(sycl_event);
@@ -313,7 +313,7 @@ ONEMKL_EXPORT sycl::event compute_backward(descriptor_type &desc, bwd<descriptor
             void *in_ptr = in;
             void *out_ptr = out;
             detail::execute_checked(func_name, plan, &in_ptr, &out_ptr, info);
-            detail::sync_checked(func_name, stream);
+            
         });
     });
     commit->set_last_usm_workspace_event_if_rqd(sycl_event);
@@ -345,7 +345,7 @@ ONEMKL_EXPORT sycl::event compute_backward(descriptor_type &desc, scalar<descrip
             std::array<void *, 2> in_native{ in_re + offsets[0], in_im + offsets[0] };
             std::array<void *, 2> out_native{ out_re + offsets[1], out_im + offsets[1] };
             detail::execute_checked(func_name, plan, in_native.data(), out_native.data(), info);
-            detail::sync_checked(func_name, stream);
+            
         });
     });
     commit->set_last_usm_workspace_event_if_rqd(sycl_event);
diff --git a/src/dft/backends/rocfft/execute_helper.hpp b/src/dft/backends/rocfft/execute_helper.hpp
index 4dff6831d..8deb18d1f 100644
--- a/src/dft/backends/rocfft/execute_helper.hpp
+++ b/src/dft/backends/rocfft/execute_helper.hpp
@@ -75,21 +75,17 @@ inline hipStream_t setup_stream(const std::string &func, sycl::interop_handle &i
     return stream;
 }
 
-inline void sync_checked(const std::string &func, hipStream_t stream) {
-    auto result = hipStreamSynchronize(stream);
-    if (result != hipSuccess) {
-        throw oneapi::mkl::exception("dft/backends/rocfft", func,
-                                     "hipStreamSynchronize returned " + std::to_string(result));
-    }
-}
+//inline void sync_checked(const std::string &func, hipStream_t stream) {
+//    auto result = hipStreamSynchronize(stream);
+//    if (result != hipSuccess) {
+//        throw oneapi::mkl::exception("dft/backends/rocfft", func,
+//                                     "hipStreamSynchronize returned " + std::to_string(result));
+//    }
+//}
 
 inline void execute_checked(const std::string &func, const rocfft_plan plan, void *in_buffer[],
                             void *out_buffer[], rocfft_execution_info info) {
-    auto result = rocfft_execute(plan, in_buffer, out_buffer, info);
-    if (result != rocfft_status_success) {
-        throw oneapi::mkl::exception("dft/backends/rocfft", func,
-                                     "rocfft_execute returned " + std::to_string(result));
-    }
+    rocfft_execute(plan, in_buffer, out_buffer, info);
 }
 
 } // namespace oneapi::mkl::dft::rocfft::detail
diff --git a/src/dft/backends/rocfft/forward.cpp b/src/dft/backends/rocfft/forward.cpp
index fc3fe5900..0461542b0 100644
--- a/src/dft/backends/rocfft/forward.cpp
+++ b/src/dft/backends/rocfft/forward.cpp
@@ -94,7 +94,7 @@ ONEMKL_EXPORT void compute_forward(descriptor_type &desc,
                 reinterpret_cast<fwd<descriptor_type> *>(detail::native_mem(ih, inout_acc)) +
                 offsets[0]);
             detail::execute_checked(func_name, plan, &inout_native, nullptr, info);
-            detail::sync_checked(func_name, stream);
+            
         });
     });
 }
@@ -134,7 +134,7 @@ ONEMKL_EXPORT void compute_forward(descriptor_type &desc,
                                          offsets[0])
             };
             detail::execute_checked(func_name, plan, inout_native.data(), nullptr, info);
-            detail::sync_checked(func_name, stream);
+            
         });
     });
 }
@@ -167,7 +167,7 @@ ONEMKL_EXPORT void compute_forward(descriptor_type &desc, sycl::buffer<fwd<descr
                 reinterpret_cast<bwd<descriptor_type> *>(detail::native_mem(ih, out_acc)) +
                 offsets[1]);
             detail::execute_checked(func_name, plan, &in_native, &out_native, info);
-            detail::sync_checked(func_name, stream);
+            
         });
     });
 }
@@ -213,7 +213,7 @@ ONEMKL_EXPORT void compute_forward(descriptor_type &desc,
                                          offsets[1])
             };
             detail::execute_checked(func_name, plan, in_native.data(), out_native.data(), info);
-            detail::sync_checked(func_name, stream);
+            
         });
     });
 }
@@ -252,7 +252,7 @@ ONEMKL_EXPORT sycl::event compute_forward(descriptor_type &desc, fwd<descriptor_
 
             void *inout_ptr = inout;
             detail::execute_checked(func_name, plan, &inout_ptr, nullptr, info);
-            detail::sync_checked(func_name, stream);
+            
         });
     });
     commit->set_last_usm_workspace_event_if_rqd(sycl_event);
@@ -285,7 +285,7 @@ ONEMKL_EXPORT sycl::event compute_forward(descriptor_type &desc, scalar<descript
 
             std::array<void *, 2> inout_native{ inout_re + offsets[0], inout_im + offsets[0] };
             detail::execute_checked(func_name, plan, inout_native.data(), nullptr, info);
-            detail::sync_checked(func_name, stream);
+            
         });
     });
     commit->set_last_usm_workspace_event_if_rqd(sycl_event);
@@ -319,7 +319,7 @@ ONEMKL_EXPORT sycl::event compute_forward(descriptor_type &desc, fwd<descriptor_
             void *in_ptr = in;
             void *out_ptr = out;
             detail::execute_checked(func_name, plan, &in_ptr, &out_ptr, info);
-            detail::sync_checked(func_name, stream);
+            
         });
     });
     commit->set_last_usm_workspace_event_if_rqd(sycl_event);
@@ -351,7 +351,7 @@ ONEMKL_EXPORT sycl::event compute_forward(descriptor_type &desc, scalar<descript
             std::array<void *, 2> in_native{ in_re + offsets[0], in_im + offsets[0] };
             std::array<void *, 2> out_native{ out_re + offsets[1], out_im + offsets[1] };
             detail::execute_checked(func_name, plan, in_native.data(), out_native.data(), info);
-            detail::sync_checked(func_name, stream);
+            
         });
     });
     commit->set_last_usm_workspace_event_if_rqd(sycl_event);

From 75dd616cd780f972b520b97bef7cf1637f153dc1 Mon Sep 17 00:00:00 2001
From: Rafal Bielski <rafal.bielski@codeplay.com>
Date: Tue, 18 Jun 2024 10:26:55 +0100
Subject: [PATCH 03/10] Use AdaptiveCpp_enqueue_custom_operation in cuFFT

---
 src/dft/backends/cufft/backward.cpp       |  9 +++++----
 src/dft/backends/cufft/execute_helper.hpp | 10 +++++-----
 src/dft/backends/cufft/forward.cpp        |  9 +++++----
 3 files changed, 15 insertions(+), 13 deletions(-)

diff --git a/src/dft/backends/cufft/backward.cpp b/src/dft/backends/cufft/backward.cpp
index aea9f232f..191c76a99 100644
--- a/src/dft/backends/cufft/backward.cpp
+++ b/src/dft/backends/cufft/backward.cpp
@@ -18,6 +18,7 @@
 *******************************************************************************/
 
 #if __has_include(<sycl/sycl.hpp>)
+#include <sycl/ext/adaptivecpp/custom_operation.hpp>
 #include <sycl/sycl.hpp>
 #else
 #include <CL/sycl.hpp>
@@ -71,7 +72,7 @@ ONEMKL_EXPORT void compute_backward(descriptor_type &desc,
         auto inout_acc = inout.template get_access<sycl::access::mode::read_write>(cgh);
         commit->add_buffer_workspace_dependency_if_rqd("compute_backward", cgh);
 
-        cgh.host_task([=](sycl::interop_handle ih) {
+        cgh.AdaptiveCpp_enqueue_custom_operation([=](sycl::interop_handle ih) {
             auto stream = detail::setup_stream(func_name, ih, plan);
 
             auto inout_native = reinterpret_cast<fwd<descriptor_type> *>(
@@ -117,7 +118,7 @@ ONEMKL_EXPORT void compute_backward(descriptor_type &desc,
         auto out_acc = out.template get_access<sycl::access::mode::read_write>(cgh);
         commit->add_buffer_workspace_dependency_if_rqd("compute_backward", cgh);
 
-        cgh.host_task([=](sycl::interop_handle ih) {
+        cgh.AdaptiveCpp_enqueue_custom_operation([=](sycl::interop_handle ih) {
             auto stream = detail::setup_stream(func_name, ih, plan);
 
             auto in_native = reinterpret_cast<void *>(
@@ -171,7 +172,7 @@ ONEMKL_EXPORT sycl::event compute_backward(descriptor_type &desc, fwd<descriptor
         cgh.depends_on(dependencies);
         commit->depend_on_last_usm_workspace_event_if_rqd(cgh);
 
-        cgh.host_task([=](sycl::interop_handle ih) {
+        cgh.AdaptiveCpp_enqueue_custom_operation([=](sycl::interop_handle ih) {
             auto stream = detail::setup_stream(func_name, ih, plan);
 
             detail::cufft_execute<detail::Direction::Backward, fwd<descriptor_type>>(
@@ -217,7 +218,7 @@ ONEMKL_EXPORT sycl::event compute_backward(descriptor_type &desc, bwd<descriptor
         cgh.depends_on(dependencies);
         commit->depend_on_last_usm_workspace_event_if_rqd(cgh);
 
-        cgh.host_task([=](sycl::interop_handle ih) {
+        cgh.AdaptiveCpp_enqueue_custom_operation([=](sycl::interop_handle ih) {
             auto stream = detail::setup_stream(func_name, ih, plan);
 
             detail::cufft_execute<detail::Direction::Backward, fwd<descriptor_type>>(
diff --git a/src/dft/backends/cufft/execute_helper.hpp b/src/dft/backends/cufft/execute_helper.hpp
index 776f0f254..bdb9ec885 100644
--- a/src/dft/backends/cufft/execute_helper.hpp
+++ b/src/dft/backends/cufft/execute_helper.hpp
@@ -126,11 +126,11 @@ void cufft_execute(const std::string &func, CUstream stream, cufftHandle plan, v
         }
     }
 
-    auto result = cuStreamSynchronize(stream);
-    if (result != CUDA_SUCCESS) {
-        throw oneapi::mkl::exception("dft/backends/cufft", func,
-                                     "cuStreamSynchronize returned " + std::to_string(result));
-    }
+    // auto result = cuStreamSynchronize(stream);
+    // if (result != CUDA_SUCCESS) {
+    //     throw oneapi::mkl::exception("dft/backends/cufft", func,
+    //                                  "cuStreamSynchronize returned " + std::to_string(result));
+    // }
 }
 
 inline CUstream setup_stream(const std::string &func, sycl::interop_handle ih, cufftHandle plan) {
diff --git a/src/dft/backends/cufft/forward.cpp b/src/dft/backends/cufft/forward.cpp
index fb323c085..4bd7584af 100644
--- a/src/dft/backends/cufft/forward.cpp
+++ b/src/dft/backends/cufft/forward.cpp
@@ -19,6 +19,7 @@
 
 #include <type_traits>
 #if __has_include(<sycl/sycl.hpp>)
+#include <sycl/ext/adaptivecpp/custom_operation.hpp>
 #include <sycl/sycl.hpp>
 #else
 #include <CL/sycl.hpp>
@@ -74,7 +75,7 @@ ONEMKL_EXPORT void compute_forward(descriptor_type &desc,
         auto inout_acc = inout.template get_access<sycl::access::mode::read_write>(cgh);
         commit->add_buffer_workspace_dependency_if_rqd("compute_forward", cgh);
 
-        cgh.host_task([=](sycl::interop_handle ih) {
+        cgh.AdaptiveCpp_enqueue_custom_operation([=](sycl::interop_handle ih) {
             auto stream = detail::setup_stream(func_name, ih, plan);
 
             auto inout_native = reinterpret_cast<fwd<descriptor_type> *>(
@@ -119,7 +120,7 @@ ONEMKL_EXPORT void compute_forward(descriptor_type &desc, sycl::buffer<fwd<descr
         auto out_acc = out.template get_access<sycl::access::mode::read_write>(cgh);
         commit->add_buffer_workspace_dependency_if_rqd("compute_forward", cgh);
 
-        cgh.host_task([=](sycl::interop_handle ih) {
+        cgh.AdaptiveCpp_enqueue_custom_operation([=](sycl::interop_handle ih) {
             auto stream = detail::setup_stream(func_name, ih, plan);
 
             auto in_native = reinterpret_cast<void *>(
@@ -173,7 +174,7 @@ ONEMKL_EXPORT sycl::event compute_forward(descriptor_type &desc, fwd<descriptor_
         cgh.depends_on(dependencies);
         commit->depend_on_last_usm_workspace_event_if_rqd(cgh);
 
-        cgh.host_task([=](sycl::interop_handle ih) {
+        cgh.AdaptiveCpp_enqueue_custom_operation([=](sycl::interop_handle ih) {
             auto stream = detail::setup_stream(func_name, ih, plan);
 
             detail::cufft_execute<detail::Direction::Forward, fwd<descriptor_type>>(
@@ -219,7 +220,7 @@ ONEMKL_EXPORT sycl::event compute_forward(descriptor_type &desc, fwd<descriptor_
         cgh.depends_on(dependencies);
         commit->depend_on_last_usm_workspace_event_if_rqd(cgh);
 
-        cgh.host_task([=](sycl::interop_handle ih) {
+        cgh.AdaptiveCpp_enqueue_custom_operation([=](sycl::interop_handle ih) {
             auto stream = detail::setup_stream(func_name, ih, plan);
 
             detail::cufft_execute<detail::Direction::Forward, fwd<descriptor_type>>(

From b359285d6ab4389881ac2bc5aec432b3400134ec Mon Sep 17 00:00:00 2001
From: Hugh Bird <hugh.bird@codeplay.com>
Date: Tue, 2 Jul 2024 10:29:45 +0100
Subject: [PATCH 04/10] Updated API (API still having minor changes)

---
 src/dft/backends/cufft/backward.cpp  | 2 +-
 src/dft/backends/cufft/forward.cpp   | 2 +-
 src/dft/backends/rocfft/backward.cpp | 1 -
 src/dft/backends/rocfft/forward.cpp  | 1 -
 4 files changed, 2 insertions(+), 4 deletions(-)

diff --git a/src/dft/backends/cufft/backward.cpp b/src/dft/backends/cufft/backward.cpp
index 191c76a99..db19e5c06 100644
--- a/src/dft/backends/cufft/backward.cpp
+++ b/src/dft/backends/cufft/backward.cpp
@@ -18,7 +18,7 @@
 *******************************************************************************/
 
 #if __has_include(<sycl/sycl.hpp>)
-#include <sycl/ext/adaptivecpp/custom_operation.hpp>
+#define SYCL_EXT_ACPP_ENQUEUE_CUSTOM_OPERATION 1
 #include <sycl/sycl.hpp>
 #else
 #include <CL/sycl.hpp>
diff --git a/src/dft/backends/cufft/forward.cpp b/src/dft/backends/cufft/forward.cpp
index 4bd7584af..1a6887cc1 100644
--- a/src/dft/backends/cufft/forward.cpp
+++ b/src/dft/backends/cufft/forward.cpp
@@ -18,8 +18,8 @@
 *******************************************************************************/
 
 #include <type_traits>
+#define SYCL_EXT_ACPP_ENQUEUE_CUSTOM_OPERATION 1
 #if __has_include(<sycl/sycl.hpp>)
-#include <sycl/ext/adaptivecpp/custom_operation.hpp>
 #include <sycl/sycl.hpp>
 #else
 #include <CL/sycl.hpp>
diff --git a/src/dft/backends/rocfft/backward.cpp b/src/dft/backends/rocfft/backward.cpp
index 46b44d724..5c2ebf4d8 100644
--- a/src/dft/backends/rocfft/backward.cpp
+++ b/src/dft/backends/rocfft/backward.cpp
@@ -17,7 +17,6 @@
 * SPDX-License-Identifier: Apache-2.0
 *******************************************************************************/
 
-#include <sycl/ext/adaptivecpp/custom_operation.hpp>
 #if __has_include(<sycl/sycl.hpp>)
 #include <sycl/sycl.hpp>
 #else
diff --git a/src/dft/backends/rocfft/forward.cpp b/src/dft/backends/rocfft/forward.cpp
index 0461542b0..897655450 100644
--- a/src/dft/backends/rocfft/forward.cpp
+++ b/src/dft/backends/rocfft/forward.cpp
@@ -19,7 +19,6 @@
 
 #include <type_traits>
 
-#include <sycl/ext/adaptivecpp/custom_operation.hpp>
 #if __has_include(<sycl/sycl.hpp>)
 #include <sycl/sycl.hpp>
 #else

From dcef3b7875c75e115a2042d6c88f7939ac340ad9 Mon Sep 17 00:00:00 2001
From: Hugh Bird <hugh.bird@codeplay.com>
Date: Fri, 16 Aug 2024 19:42:51 +0100
Subject: [PATCH 05/10] Update to final API; Enable based on feature macro

---
 src/dft/backends/cufft/backward.cpp        |  9 ++--
 src/dft/backends/cufft/execute_helper.hpp  | 39 +++++++++++++++---
 src/dft/backends/cufft/forward.cpp         |  9 ++--
 src/dft/backends/rocfft/backward.cpp       | 39 ++++++++----------
 src/dft/backends/rocfft/execute_helper.hpp | 48 ++++++++++++++++++----
 src/dft/backends/rocfft/forward.cpp        | 44 ++++++++------------
 6 files changed, 112 insertions(+), 76 deletions(-)

diff --git a/src/dft/backends/cufft/backward.cpp b/src/dft/backends/cufft/backward.cpp
index db19e5c06..693ad4d1b 100644
--- a/src/dft/backends/cufft/backward.cpp
+++ b/src/dft/backends/cufft/backward.cpp
@@ -18,7 +18,6 @@
 *******************************************************************************/
 
 #if __has_include(<sycl/sycl.hpp>)
-#define SYCL_EXT_ACPP_ENQUEUE_CUSTOM_OPERATION 1
 #include <sycl/sycl.hpp>
 #else
 #include <CL/sycl.hpp>
@@ -72,7 +71,7 @@ ONEMKL_EXPORT void compute_backward(descriptor_type &desc,
         auto inout_acc = inout.template get_access<sycl::access::mode::read_write>(cgh);
         commit->add_buffer_workspace_dependency_if_rqd("compute_backward", cgh);
 
-        cgh.AdaptiveCpp_enqueue_custom_operation([=](sycl::interop_handle ih) {
+        detail::cufft_enqueue_task(cgh, [=](sycl::interop_handle ih) {
             auto stream = detail::setup_stream(func_name, ih, plan);
 
             auto inout_native = reinterpret_cast<fwd<descriptor_type> *>(
@@ -118,7 +117,7 @@ ONEMKL_EXPORT void compute_backward(descriptor_type &desc,
         auto out_acc = out.template get_access<sycl::access::mode::read_write>(cgh);
         commit->add_buffer_workspace_dependency_if_rqd("compute_backward", cgh);
 
-        cgh.AdaptiveCpp_enqueue_custom_operation([=](sycl::interop_handle ih) {
+        detail::cufft_enqueue_task(cgh, [=](sycl::interop_handle ih) {
             auto stream = detail::setup_stream(func_name, ih, plan);
 
             auto in_native = reinterpret_cast<void *>(
@@ -172,7 +171,7 @@ ONEMKL_EXPORT sycl::event compute_backward(descriptor_type &desc, fwd<descriptor
         cgh.depends_on(dependencies);
         commit->depend_on_last_usm_workspace_event_if_rqd(cgh);
 
-        cgh.AdaptiveCpp_enqueue_custom_operation([=](sycl::interop_handle ih) {
+        detail::cufft_enqueue_task(cgh, [=](sycl::interop_handle ih) {
             auto stream = detail::setup_stream(func_name, ih, plan);
 
             detail::cufft_execute<detail::Direction::Backward, fwd<descriptor_type>>(
@@ -218,7 +217,7 @@ ONEMKL_EXPORT sycl::event compute_backward(descriptor_type &desc, bwd<descriptor
         cgh.depends_on(dependencies);
         commit->depend_on_last_usm_workspace_event_if_rqd(cgh);
 
-        cgh.AdaptiveCpp_enqueue_custom_operation([=](sycl::interop_handle ih) {
+        detail::cufft_enqueue_task(cgh, [=](sycl::interop_handle ih) {
             auto stream = detail::setup_stream(func_name, ih, plan);
 
             detail::cufft_execute<detail::Direction::Backward, fwd<descriptor_type>>(
diff --git a/src/dft/backends/cufft/execute_helper.hpp b/src/dft/backends/cufft/execute_helper.hpp
index bdb9ec885..f33c6eeca 100644
--- a/src/dft/backends/cufft/execute_helper.hpp
+++ b/src/dft/backends/cufft/execute_helper.hpp
@@ -20,6 +20,8 @@
 #ifndef _ONEMKL_DFT_SRC_CUFFT_EXECUTE_HPP_
 #define _ONEMKL_DFT_SRC_CUFFT_EXECUTE_HPP_
 
+#define SYCL_EXT_ACPP_ENQUEUE_CUSTOM_OPERATION
+
 #if __has_include(<sycl/sycl.hpp>)
 #include <sycl/sycl.hpp>
 #else
@@ -125,12 +127,16 @@ void cufft_execute(const std::string &func, CUstream stream, cufftHandle plan, v
             }
         }
     }
-
-    // auto result = cuStreamSynchronize(stream);
-    // if (result != CUDA_SUCCESS) {
-    //     throw oneapi::mkl::exception("dft/backends/cufft", func,
-    //                                  "cuStreamSynchronize returned " + std::to_string(result));
-    // }
+#ifndef SYCL_EXT_ONEAPI_ENQUEUE_NATIVE_COMMAND
+    // If not using the enqueue native extension, the host task must wait on the
+    // asynchronous operation to complete. Otherwise it report the operation
+    // as complete early.
+    auto result = cuStreamSynchronize(stream);
+    if (result != CUDA_SUCCESS) {
+        throw oneapi::mkl::exception("dft/backends/cufft", func,
+                                     "cuStreamSynchronize returned " + std::to_string(result));
+    }
+#endif
 }
 
 inline CUstream setup_stream(const std::string &func, sycl::interop_handle ih, cufftHandle plan) {
@@ -143,6 +149,27 @@ inline CUstream setup_stream(const std::string &func, sycl::interop_handle ih, c
     return stream;
 }
 
+
+/** Wrap interop API to launch interop host task.
+ * 
+ * @tparam HandlerT The command group handler type
+ * @tparam FnT The body of the enqueued task
+ *
+ * Either uses host task interop API, or enqueue native command extension.
+ * This extension avoids host synchronization after 
+ * the CUDA call is complete.
+ */
+template <typename HandlerT, typename FnT>
+static inline void cufft_enqueue_task(HandlerT&& cgh, FnT&& f) {
+#ifdef SYCL_EXT_ONEAPI_ENQUEUE_NATIVE_COMMAND
+    cgh.ext_codeplay_enqueue_native_command([=](sycl::interop_handle ih){
+#else
+    cgh.host_task([=](sycl::interop_handle ih){
+#endif
+        f(std::move(ih));
+    });
+}
+
 } // namespace oneapi::mkl::dft::cufft::detail
 
 #endif
diff --git a/src/dft/backends/cufft/forward.cpp b/src/dft/backends/cufft/forward.cpp
index 1a6887cc1..bdbda2cb5 100644
--- a/src/dft/backends/cufft/forward.cpp
+++ b/src/dft/backends/cufft/forward.cpp
@@ -18,7 +18,6 @@
 *******************************************************************************/
 
 #include <type_traits>
-#define SYCL_EXT_ACPP_ENQUEUE_CUSTOM_OPERATION 1
 #if __has_include(<sycl/sycl.hpp>)
 #include <sycl/sycl.hpp>
 #else
@@ -75,7 +74,7 @@ ONEMKL_EXPORT void compute_forward(descriptor_type &desc,
         auto inout_acc = inout.template get_access<sycl::access::mode::read_write>(cgh);
         commit->add_buffer_workspace_dependency_if_rqd("compute_forward", cgh);
 
-        cgh.AdaptiveCpp_enqueue_custom_operation([=](sycl::interop_handle ih) {
+        detail::cufft_enqueue_task(cgh, [=](sycl::interop_handle ih) {
             auto stream = detail::setup_stream(func_name, ih, plan);
 
             auto inout_native = reinterpret_cast<fwd<descriptor_type> *>(
@@ -120,7 +119,7 @@ ONEMKL_EXPORT void compute_forward(descriptor_type &desc, sycl::buffer<fwd<descr
         auto out_acc = out.template get_access<sycl::access::mode::read_write>(cgh);
         commit->add_buffer_workspace_dependency_if_rqd("compute_forward", cgh);
 
-        cgh.AdaptiveCpp_enqueue_custom_operation([=](sycl::interop_handle ih) {
+        detail::cufft_enqueue_task(cgh, [=](sycl::interop_handle ih) {
             auto stream = detail::setup_stream(func_name, ih, plan);
 
             auto in_native = reinterpret_cast<void *>(
@@ -174,7 +173,7 @@ ONEMKL_EXPORT sycl::event compute_forward(descriptor_type &desc, fwd<descriptor_
         cgh.depends_on(dependencies);
         commit->depend_on_last_usm_workspace_event_if_rqd(cgh);
 
-        cgh.AdaptiveCpp_enqueue_custom_operation([=](sycl::interop_handle ih) {
+        detail::cufft_enqueue_task(cgh, [=](sycl::interop_handle ih) {
             auto stream = detail::setup_stream(func_name, ih, plan);
 
             detail::cufft_execute<detail::Direction::Forward, fwd<descriptor_type>>(
@@ -220,7 +219,7 @@ ONEMKL_EXPORT sycl::event compute_forward(descriptor_type &desc, fwd<descriptor_
         cgh.depends_on(dependencies);
         commit->depend_on_last_usm_workspace_event_if_rqd(cgh);
 
-        cgh.AdaptiveCpp_enqueue_custom_operation([=](sycl::interop_handle ih) {
+        detail::cufft_enqueue_task(cgh, [=](sycl::interop_handle ih) {
             auto stream = detail::setup_stream(func_name, ih, plan);
 
             detail::cufft_execute<detail::Direction::Forward, fwd<descriptor_type>>(
diff --git a/src/dft/backends/rocfft/backward.cpp b/src/dft/backends/rocfft/backward.cpp
index 5c2ebf4d8..bdb1c9638 100644
--- a/src/dft/backends/rocfft/backward.cpp
+++ b/src/dft/backends/rocfft/backward.cpp
@@ -78,14 +78,13 @@ ONEMKL_EXPORT void compute_backward(descriptor_type &desc,
         auto inout_acc = inout.template get_access<sycl::access::mode::read_write>(cgh);
         commit->add_buffer_workspace_dependency_if_rqd("compute_backward", cgh);
 
-        cgh.AdaptiveCpp_enqueue_custom_operation([=](sycl::interop_handle ih) {
+        detail::rocfft_enqueue_task(cgh, [=](sycl::interop_handle ih) {
             auto stream = detail::setup_stream(func_name, ih, info);
 
             auto inout_native = reinterpret_cast<void *>(
                 reinterpret_cast<fwd<descriptor_type> *>(detail::native_mem(ih, inout_acc)) +
                 offsets[0]);
-            detail::execute_checked(func_name, plan, &inout_native, nullptr, info);
-            
+            detail::execute_checked(func_name, stream,  plan, &inout_native, nullptr, info);
         });
     });
 }
@@ -113,7 +112,7 @@ ONEMKL_EXPORT void compute_backward(descriptor_type &desc,
         auto inout_im_acc = inout_im.template get_access<sycl::access::mode::read_write>(cgh);
         commit->add_buffer_workspace_dependency_if_rqd("compute_backward", cgh);
 
-        cgh.AdaptiveCpp_enqueue_custom_operation([=](sycl::interop_handle ih) {
+        detail::rocfft_enqueue_task(cgh, [=](sycl::interop_handle ih) {
             auto stream = detail::setup_stream(func_name, ih, info);
 
             std::array<void *, 2> inout_native{
@@ -124,8 +123,7 @@ ONEMKL_EXPORT void compute_backward(descriptor_type &desc,
                                              detail::native_mem(ih, inout_im_acc)) +
                                          offsets[0])
             };
-            detail::execute_checked(func_name, plan, inout_native.data(), nullptr, info);
-            
+            detail::execute_checked(func_name, stream,  plan, inout_native.data(), nullptr, info);
         });
     });
 }
@@ -148,7 +146,7 @@ ONEMKL_EXPORT void compute_backward(descriptor_type &desc,
         auto out_acc = out.template get_access<sycl::access::mode::read_write>(cgh);
         commit->add_buffer_workspace_dependency_if_rqd("compute_backward", cgh);
 
-        cgh.AdaptiveCpp_enqueue_custom_operation([=](sycl::interop_handle ih) {
+        detail::rocfft_enqueue_task(cgh, [=](sycl::interop_handle ih) {
             const std::string func_name = "compute_backward(desc, in, out)";
             auto stream = detail::setup_stream(func_name, ih, info);
 
@@ -158,8 +156,7 @@ ONEMKL_EXPORT void compute_backward(descriptor_type &desc,
             auto out_native = reinterpret_cast<void *>(
                 reinterpret_cast<fwd<descriptor_type> *>(detail::native_mem(ih, out_acc)) +
                 offsets[1]);
-            detail::execute_checked(func_name, plan, &in_native, &out_native, info);
-            
+            detail::execute_checked(func_name, stream,  plan, &in_native, &out_native, info);
         });
     });
 }
@@ -184,7 +181,7 @@ ONEMKL_EXPORT void compute_backward(descriptor_type &desc,
         auto out_im_acc = out_im.template get_access<sycl::access::mode::read_write>(cgh);
         commit->add_buffer_workspace_dependency_if_rqd("compute_backward", cgh);
 
-        cgh.AdaptiveCpp_enqueue_custom_operation([=](sycl::interop_handle ih) {
+        detail::rocfft_enqueue_task(cgh, [=](sycl::interop_handle ih) {
             const std::string func_name = "compute_backward(desc, in_re, in_im, out_re, out_im)";
             auto stream = detail::setup_stream(func_name, ih, info);
 
@@ -204,8 +201,7 @@ ONEMKL_EXPORT void compute_backward(descriptor_type &desc,
                                              detail::native_mem(ih, out_im_acc)) +
                                          offsets[1])
             };
-            detail::execute_checked(func_name, plan, in_native.data(), out_native.data(), info);
-            
+            detail::execute_checked(func_name, stream,  plan, in_native.data(), out_native.data(), info);
         });
     });
 }
@@ -239,12 +235,11 @@ ONEMKL_EXPORT sycl::event compute_backward(descriptor_type &desc, fwd<descriptor
         cgh.depends_on(deps);
         commit->depend_on_last_usm_workspace_event_if_rqd(cgh);
 
-        cgh.AdaptiveCpp_enqueue_custom_operation([=](sycl::interop_handle ih) {
+        detail::rocfft_enqueue_task(cgh, [=](sycl::interop_handle ih) {
             auto stream = detail::setup_stream(func_name, ih, info);
 
             void *inout_ptr = inout;
-            detail::execute_checked(func_name, plan, &inout_ptr, nullptr, info);
-            
+            detail::execute_checked(func_name, stream,  plan, &inout_ptr, nullptr, info);
         });
     });
     commit->set_last_usm_workspace_event_if_rqd(sycl_event);
@@ -273,11 +268,11 @@ ONEMKL_EXPORT sycl::event compute_backward(descriptor_type &desc, scalar<descrip
         cgh.depends_on(deps);
         commit->depend_on_last_usm_workspace_event_if_rqd(cgh);
 
-        cgh.AdaptiveCpp_enqueue_custom_operation([=](sycl::interop_handle ih) {
+        detail::rocfft_enqueue_task(cgh, [=](sycl::interop_handle ih) {
             auto stream = detail::setup_stream(func_name, ih, info);
 
             std::array<void *, 2> inout_native{ inout_re + offsets[0], inout_im + offsets[0] };
-            detail::execute_checked(func_name, plan, inout_native.data(), nullptr, info);
+            detail::execute_checked(func_name, stream,  plan, inout_native.data(), nullptr, info);
             
         });
     });
@@ -305,14 +300,13 @@ ONEMKL_EXPORT sycl::event compute_backward(descriptor_type &desc, bwd<descriptor
         cgh.depends_on(deps);
         commit->depend_on_last_usm_workspace_event_if_rqd(cgh);
 
-        cgh.AdaptiveCpp_enqueue_custom_operation([=](sycl::interop_handle ih) {
+        detail::rocfft_enqueue_task(cgh, [=](sycl::interop_handle ih) {
             const std::string func_name = "compute_backward(desc, in, out, deps)";
             auto stream = detail::setup_stream(func_name, ih, info);
 
             void *in_ptr = in;
             void *out_ptr = out;
-            detail::execute_checked(func_name, plan, &in_ptr, &out_ptr, info);
-            
+            detail::execute_checked(func_name, stream,  plan, &in_ptr, &out_ptr, info);
         });
     });
     commit->set_last_usm_workspace_event_if_rqd(sycl_event);
@@ -336,15 +330,14 @@ ONEMKL_EXPORT sycl::event compute_backward(descriptor_type &desc, scalar<descrip
         cgh.depends_on(deps);
         commit->depend_on_last_usm_workspace_event_if_rqd(cgh);
 
-        cgh.AdaptiveCpp_enqueue_custom_operation([=](sycl::interop_handle ih) {
+        detail::rocfft_enqueue_task(cgh, [=](sycl::interop_handle ih) {
             const std::string func_name =
                 "compute_backward(desc, in_re, in_im, out_re, out_im, deps)";
             auto stream = detail::setup_stream(func_name, ih, info);
 
             std::array<void *, 2> in_native{ in_re + offsets[0], in_im + offsets[0] };
             std::array<void *, 2> out_native{ out_re + offsets[1], out_im + offsets[1] };
-            detail::execute_checked(func_name, plan, in_native.data(), out_native.data(), info);
-            
+            detail::execute_checked(func_name, stream,  plan, in_native.data(), out_native.data(), info);
         });
     });
     commit->set_last_usm_workspace_event_if_rqd(sycl_event);
diff --git a/src/dft/backends/rocfft/execute_helper.hpp b/src/dft/backends/rocfft/execute_helper.hpp
index 8deb18d1f..168774e93 100644
--- a/src/dft/backends/rocfft/execute_helper.hpp
+++ b/src/dft/backends/rocfft/execute_helper.hpp
@@ -75,17 +75,47 @@ inline hipStream_t setup_stream(const std::string &func, sycl::interop_handle &i
     return stream;
 }
 
-//inline void sync_checked(const std::string &func, hipStream_t stream) {
-//    auto result = hipStreamSynchronize(stream);
-//    if (result != hipSuccess) {
-//        throw oneapi::mkl::exception("dft/backends/rocfft", func,
-//                                     "hipStreamSynchronize returned " + std::to_string(result));
-//    }
-//}
+inline void sync_checked(const std::string &func, hipStream_t stream) {
+   auto result = hipStreamSynchronize(stream);
+   if (result != hipSuccess) {
+       throw oneapi::mkl::exception("dft/backends/rocfft", func,
+                                    "hipStreamSynchronize returned " + std::to_string(result));
+   }
+}
 
-inline void execute_checked(const std::string &func, const rocfft_plan plan, void *in_buffer[],
+inline void execute_checked(const std::string &func, hipStream_t stream, const rocfft_plan plan, void *in_buffer[],
                             void *out_buffer[], rocfft_execution_info info) {
-    rocfft_execute(plan, in_buffer, out_buffer, info);
+    auto result = rocfft_execute(plan, in_buffer, out_buffer, info);    
+    if (result != rocfft_status_success) {
+        throw oneapi::mkl::exception("dft/backends/rocfft", func,
+                                     "rocfft_execute returned " + std::to_string(result));
+    }
+#ifndef SYCL_EXT_ONEAPI_ENQUEUE_NATIVE_COMMAND
+    // If not using equeue native extension, the host task must wait on the
+    // asynchronous operation to complete. Otherwise it report the operation
+    // as complete early.
+    sync_checked(func, stream);
+#endif
+}
+
+/** Wrap interop API to launch interop host task.
+ * 
+ * @tparam HandlerT The command group handler type
+ * @tparam FnT The body of the enqueued task
+ *
+ * Either uses host task interop API, or enqueue native command extension.
+ * This extension avoids host synchronization after 
+ * the CUDA call is complete.
+ */
+template <typename HandlerT, typename FnT>
+static inline void rocfft_enqueue_task(HandlerT&& cgh, FnT&& f) {
+#ifdef SYCL_EXT_ONEAPI_ENQUEUE_NATIVE_COMMAND
+    cgh.ext_codeplay_enqueue_native_command([=](sycl::interop_handle ih){
+#else
+    cgh.host_task([=](sycl::interop_handle ih){
+#endif
+        f(std::move(ih));
+    });
 }
 
 } // namespace oneapi::mkl::dft::rocfft::detail
diff --git a/src/dft/backends/rocfft/forward.cpp b/src/dft/backends/rocfft/forward.cpp
index 897655450..2f2212041 100644
--- a/src/dft/backends/rocfft/forward.cpp
+++ b/src/dft/backends/rocfft/forward.cpp
@@ -24,10 +24,6 @@
 #else
 #include <CL/sycl.hpp>
 #endif
-#ifndef SYCL_EXT_ACPP_ENQUEUE_CUSTOM_OPERATION
-# error Extension not implemented!
-#endif
-
 
 #include "oneapi/mkl/exceptions.hpp"
 
@@ -86,14 +82,13 @@ ONEMKL_EXPORT void compute_forward(descriptor_type &desc,
         auto inout_acc = inout.template get_access<sycl::access::mode::read_write>(cgh);
         commit->add_buffer_workspace_dependency_if_rqd("compute_forward", cgh);
 
-        cgh.AdaptiveCpp_enqueue_custom_operation([=](sycl::interop_handle ih) {
+        detail::rocfft_enqueue_task(cgh, [=](sycl::interop_handle ih) {
             auto stream = detail::setup_stream(func_name, ih, info);
 
             auto inout_native = reinterpret_cast<void *>(
                 reinterpret_cast<fwd<descriptor_type> *>(detail::native_mem(ih, inout_acc)) +
                 offsets[0]);
-            detail::execute_checked(func_name, plan, &inout_native, nullptr, info);
-            
+            detail::execute_checked(func_name, stream,  plan, &inout_native, nullptr, info);
         });
     });
 }
@@ -121,7 +116,7 @@ ONEMKL_EXPORT void compute_forward(descriptor_type &desc,
         auto inout_im_acc = inout_im.template get_access<sycl::access::mode::read_write>(cgh);
         commit->add_buffer_workspace_dependency_if_rqd("compute_forward", cgh);
 
-        cgh.AdaptiveCpp_enqueue_custom_operation([=](sycl::interop_handle ih) {
+        detail::rocfft_enqueue_task(cgh, [=](sycl::interop_handle ih) {
             auto stream = detail::setup_stream(func_name, ih, info);
 
             std::array<void *, 2> inout_native{
@@ -132,8 +127,7 @@ ONEMKL_EXPORT void compute_forward(descriptor_type &desc,
                                              detail::native_mem(ih, inout_im_acc)) +
                                          offsets[0])
             };
-            detail::execute_checked(func_name, plan, inout_native.data(), nullptr, info);
-            
+            detail::execute_checked(func_name, stream,  plan, inout_native.data(), nullptr, info);
         });
     });
 }
@@ -155,7 +149,7 @@ ONEMKL_EXPORT void compute_forward(descriptor_type &desc, sycl::buffer<fwd<descr
         auto out_acc = out.template get_access<sycl::access::mode::read_write>(cgh);
         commit->add_buffer_workspace_dependency_if_rqd("compute_forward", cgh);
 
-        cgh.AdaptiveCpp_enqueue_custom_operation([=](sycl::interop_handle ih) {
+        detail::rocfft_enqueue_task(cgh, [=](sycl::interop_handle ih) {
             const std::string func_name = "compute_forward(desc, in, out)";
             auto stream = detail::setup_stream(func_name, ih, info);
 
@@ -165,8 +159,7 @@ ONEMKL_EXPORT void compute_forward(descriptor_type &desc, sycl::buffer<fwd<descr
             auto out_native = reinterpret_cast<void *>(
                 reinterpret_cast<bwd<descriptor_type> *>(detail::native_mem(ih, out_acc)) +
                 offsets[1]);
-            detail::execute_checked(func_name, plan, &in_native, &out_native, info);
-            
+            detail::execute_checked(func_name, stream,  plan, &in_native, &out_native, info);
         });
     });
 }
@@ -191,7 +184,7 @@ ONEMKL_EXPORT void compute_forward(descriptor_type &desc,
         auto out_im_acc = out_im.template get_access<sycl::access::mode::read_write>(cgh);
         commit->add_buffer_workspace_dependency_if_rqd("compute_forward", cgh);
 
-        cgh.AdaptiveCpp_enqueue_custom_operation([=](sycl::interop_handle ih) {
+        detail::rocfft_enqueue_task(cgh, [=](sycl::interop_handle ih) {
             const std::string func_name = "compute_forward(desc, in_re, in_im, out_re, out_im)";
             auto stream = detail::setup_stream(func_name, ih, info);
 
@@ -211,8 +204,7 @@ ONEMKL_EXPORT void compute_forward(descriptor_type &desc,
                                              detail::native_mem(ih, out_im_acc)) +
                                          offsets[1])
             };
-            detail::execute_checked(func_name, plan, in_native.data(), out_native.data(), info);
-            
+            detail::execute_checked(func_name, stream,  plan, in_native.data(), out_native.data(), info);
         });
     });
 }
@@ -246,12 +238,11 @@ ONEMKL_EXPORT sycl::event compute_forward(descriptor_type &desc, fwd<descriptor_
         cgh.depends_on(deps);
         commit->depend_on_last_usm_workspace_event_if_rqd(cgh);
 
-        cgh.AdaptiveCpp_enqueue_custom_operation([=](sycl::interop_handle ih) {
+        detail::rocfft_enqueue_task(cgh, [=](sycl::interop_handle ih) {
             auto stream = detail::setup_stream(func_name, ih, info);
 
             void *inout_ptr = inout;
-            detail::execute_checked(func_name, plan, &inout_ptr, nullptr, info);
-            
+            detail::execute_checked(func_name, stream,  plan, &inout_ptr, nullptr, info);
         });
     });
     commit->set_last_usm_workspace_event_if_rqd(sycl_event);
@@ -279,12 +270,11 @@ ONEMKL_EXPORT sycl::event compute_forward(descriptor_type &desc, scalar<descript
     sycl::event sycl_event = queue.submit([&](sycl::handler &cgh) {
         cgh.depends_on(deps);
         commit->depend_on_last_usm_workspace_event_if_rqd(cgh);
-        cgh.AdaptiveCpp_enqueue_custom_operation([=](sycl::interop_handle ih) {
+        detail::rocfft_enqueue_task(cgh, [=](sycl::interop_handle ih) {
             auto stream = detail::setup_stream(func_name, ih, info);
 
             std::array<void *, 2> inout_native{ inout_re + offsets[0], inout_im + offsets[0] };
-            detail::execute_checked(func_name, plan, inout_native.data(), nullptr, info);
-            
+            detail::execute_checked(func_name, stream,  plan, inout_native.data(), nullptr, info);
         });
     });
     commit->set_last_usm_workspace_event_if_rqd(sycl_event);
@@ -311,14 +301,13 @@ ONEMKL_EXPORT sycl::event compute_forward(descriptor_type &desc, fwd<descriptor_
         cgh.depends_on(deps);
         commit->depend_on_last_usm_workspace_event_if_rqd(cgh);
 
-        cgh.AdaptiveCpp_enqueue_custom_operation([=](sycl::interop_handle ih) {
+        detail::rocfft_enqueue_task(cgh, [=](sycl::interop_handle ih) {
             const std::string func_name = "compute_forward(desc, in, out, deps)";
             auto stream = detail::setup_stream(func_name, ih, info);
 
             void *in_ptr = in;
             void *out_ptr = out;
-            detail::execute_checked(func_name, plan, &in_ptr, &out_ptr, info);
-            
+            detail::execute_checked(func_name, stream,  plan, &in_ptr, &out_ptr, info);
         });
     });
     commit->set_last_usm_workspace_event_if_rqd(sycl_event);
@@ -342,15 +331,14 @@ ONEMKL_EXPORT sycl::event compute_forward(descriptor_type &desc, scalar<descript
         cgh.depends_on(deps);
         commit->depend_on_last_usm_workspace_event_if_rqd(cgh);
 
-        cgh.AdaptiveCpp_enqueue_custom_operation([=](sycl::interop_handle ih) {
+        detail::rocfft_enqueue_task(cgh, [=](sycl::interop_handle ih) {
             const std::string func_name =
                 "compute_forward(desc, in_re, in_im, out_re, out_im, deps)";
             auto stream = detail::setup_stream(func_name, ih, info);
 
             std::array<void *, 2> in_native{ in_re + offsets[0], in_im + offsets[0] };
             std::array<void *, 2> out_native{ out_re + offsets[1], out_im + offsets[1] };
-            detail::execute_checked(func_name, plan, in_native.data(), out_native.data(), info);
-            
+            detail::execute_checked(func_name, stream,  plan, in_native.data(), out_native.data(), info);
         });
     });
     commit->set_last_usm_workspace_event_if_rqd(sycl_event);

From cdb496cfc710f91e81af2c4aad704a999278e849 Mon Sep 17 00:00:00 2001
From: JackAKirk <jack.kirk@codeplay.com>
Date: Fri, 27 Sep 2024 07:58:11 -0700
Subject: [PATCH 06/10] Remove unused macro

Signed-off-by: JackAKirk <jack.kirk@codeplay.com>
---
 src/dft/backends/cufft/execute_helper.hpp | 2 --
 src/dft/backends/rocfft/forward.cpp       | 1 -
 2 files changed, 3 deletions(-)

diff --git a/src/dft/backends/cufft/execute_helper.hpp b/src/dft/backends/cufft/execute_helper.hpp
index f33c6eeca..644cf7148 100644
--- a/src/dft/backends/cufft/execute_helper.hpp
+++ b/src/dft/backends/cufft/execute_helper.hpp
@@ -20,8 +20,6 @@
 #ifndef _ONEMKL_DFT_SRC_CUFFT_EXECUTE_HPP_
 #define _ONEMKL_DFT_SRC_CUFFT_EXECUTE_HPP_
 
-#define SYCL_EXT_ACPP_ENQUEUE_CUSTOM_OPERATION
-
 #if __has_include(<sycl/sycl.hpp>)
 #include <sycl/sycl.hpp>
 #else
diff --git a/src/dft/backends/rocfft/forward.cpp b/src/dft/backends/rocfft/forward.cpp
index 2f2212041..daacc685d 100644
--- a/src/dft/backends/rocfft/forward.cpp
+++ b/src/dft/backends/rocfft/forward.cpp
@@ -18,7 +18,6 @@
 *******************************************************************************/
 
 #include <type_traits>
-
 #if __has_include(<sycl/sycl.hpp>)
 #include <sycl/sycl.hpp>
 #else

From f15983cf555b61e616d4c6bae5935e123f3550d2 Mon Sep 17 00:00:00 2001
From: JackAKirk <jack.kirk@codeplay.com>
Date: Thu, 3 Oct 2024 14:29:47 +0100
Subject: [PATCH 07/10] Update src/dft/backends/rocfft/execute_helper.hpp

remove whitespace

Co-authored-by: Romain Biessy <romain.biessy@codeplay.com>
---
 src/dft/backends/rocfft/execute_helper.hpp | 2 +-
 1 file changed, 1 insertion(+), 1 deletion(-)

diff --git a/src/dft/backends/rocfft/execute_helper.hpp b/src/dft/backends/rocfft/execute_helper.hpp
index 168774e93..49c499637 100644
--- a/src/dft/backends/rocfft/execute_helper.hpp
+++ b/src/dft/backends/rocfft/execute_helper.hpp
@@ -85,7 +85,7 @@ inline void sync_checked(const std::string &func, hipStream_t stream) {
 
 inline void execute_checked(const std::string &func, hipStream_t stream, const rocfft_plan plan, void *in_buffer[],
                             void *out_buffer[], rocfft_execution_info info) {
-    auto result = rocfft_execute(plan, in_buffer, out_buffer, info);    
+    auto result = rocfft_execute(plan, in_buffer, out_buffer, info);
     if (result != rocfft_status_success) {
         throw oneapi::mkl::exception("dft/backends/rocfft", func,
                                      "rocfft_execute returned " + std::to_string(result));

From 1df1cb1ae35d0a608d80dd7e78c1092af00d191a Mon Sep 17 00:00:00 2001
From: JackAKirk <jack.kirk@codeplay.com>
Date: Fri, 11 Oct 2024 07:38:55 -0700
Subject: [PATCH 08/10] switch to portable fft_enqueue_task

Signed-off-by: JackAKirk <jack.kirk@codeplay.com>
---
 src/dft/backends/cufft/backward.cpp        |  9 ++--
 src/dft/backends/cufft/execute_helper.hpp  | 21 ---------
 src/dft/backends/cufft/forward.cpp         |  9 ++--
 src/dft/backends/rocfft/backward.cpp       | 17 +++----
 src/dft/backends/rocfft/execute_helper.hpp | 20 --------
 src/dft/backends/rocfft/forward.cpp        | 17 +++----
 src/dft/execute_helper_generic.hpp         | 53 ++++++++++++++++++++++
 7 files changed, 81 insertions(+), 65 deletions(-)
 create mode 100644 src/dft/execute_helper_generic.hpp

diff --git a/src/dft/backends/cufft/backward.cpp b/src/dft/backends/cufft/backward.cpp
index 693ad4d1b..80e475991 100644
--- a/src/dft/backends/cufft/backward.cpp
+++ b/src/dft/backends/cufft/backward.cpp
@@ -30,6 +30,7 @@
 #include "oneapi/mkl/dft/types.hpp"
 
 #include "execute_helper.hpp"
+#include "../../execute_helper_generic.hpp"
 
 #include <cufft.h>
 
@@ -71,7 +72,7 @@ ONEMKL_EXPORT void compute_backward(descriptor_type &desc,
         auto inout_acc = inout.template get_access<sycl::access::mode::read_write>(cgh);
         commit->add_buffer_workspace_dependency_if_rqd("compute_backward", cgh);
 
-        detail::cufft_enqueue_task(cgh, [=](sycl::interop_handle ih) {
+        dft::detail::fft_enqueue_task(cgh, [=](sycl::interop_handle ih) {
             auto stream = detail::setup_stream(func_name, ih, plan);
 
             auto inout_native = reinterpret_cast<fwd<descriptor_type> *>(
@@ -117,7 +118,7 @@ ONEMKL_EXPORT void compute_backward(descriptor_type &desc,
         auto out_acc = out.template get_access<sycl::access::mode::read_write>(cgh);
         commit->add_buffer_workspace_dependency_if_rqd("compute_backward", cgh);
 
-        detail::cufft_enqueue_task(cgh, [=](sycl::interop_handle ih) {
+        dft::detail::fft_enqueue_task(cgh, [=](sycl::interop_handle ih) {
             auto stream = detail::setup_stream(func_name, ih, plan);
 
             auto in_native = reinterpret_cast<void *>(
@@ -171,7 +172,7 @@ ONEMKL_EXPORT sycl::event compute_backward(descriptor_type &desc, fwd<descriptor
         cgh.depends_on(dependencies);
         commit->depend_on_last_usm_workspace_event_if_rqd(cgh);
 
-        detail::cufft_enqueue_task(cgh, [=](sycl::interop_handle ih) {
+        dft::detail::fft_enqueue_task(cgh, [=](sycl::interop_handle ih) {
             auto stream = detail::setup_stream(func_name, ih, plan);
 
             detail::cufft_execute<detail::Direction::Backward, fwd<descriptor_type>>(
@@ -217,7 +218,7 @@ ONEMKL_EXPORT sycl::event compute_backward(descriptor_type &desc, bwd<descriptor
         cgh.depends_on(dependencies);
         commit->depend_on_last_usm_workspace_event_if_rqd(cgh);
 
-        detail::cufft_enqueue_task(cgh, [=](sycl::interop_handle ih) {
+        dft::detail::fft_enqueue_task(cgh, [=](sycl::interop_handle ih) {
             auto stream = detail::setup_stream(func_name, ih, plan);
 
             detail::cufft_execute<detail::Direction::Backward, fwd<descriptor_type>>(
diff --git a/src/dft/backends/cufft/execute_helper.hpp b/src/dft/backends/cufft/execute_helper.hpp
index 644cf7148..bbe32c146 100644
--- a/src/dft/backends/cufft/execute_helper.hpp
+++ b/src/dft/backends/cufft/execute_helper.hpp
@@ -147,27 +147,6 @@ inline CUstream setup_stream(const std::string &func, sycl::interop_handle ih, c
     return stream;
 }
 
-
-/** Wrap interop API to launch interop host task.
- * 
- * @tparam HandlerT The command group handler type
- * @tparam FnT The body of the enqueued task
- *
- * Either uses host task interop API, or enqueue native command extension.
- * This extension avoids host synchronization after 
- * the CUDA call is complete.
- */
-template <typename HandlerT, typename FnT>
-static inline void cufft_enqueue_task(HandlerT&& cgh, FnT&& f) {
-#ifdef SYCL_EXT_ONEAPI_ENQUEUE_NATIVE_COMMAND
-    cgh.ext_codeplay_enqueue_native_command([=](sycl::interop_handle ih){
-#else
-    cgh.host_task([=](sycl::interop_handle ih){
-#endif
-        f(std::move(ih));
-    });
-}
-
 } // namespace oneapi::mkl::dft::cufft::detail
 
 #endif
diff --git a/src/dft/backends/cufft/forward.cpp b/src/dft/backends/cufft/forward.cpp
index bdbda2cb5..7cf73976d 100644
--- a/src/dft/backends/cufft/forward.cpp
+++ b/src/dft/backends/cufft/forward.cpp
@@ -31,6 +31,7 @@
 #include "oneapi/mkl/dft/types.hpp"
 
 #include "execute_helper.hpp"
+#include "../../execute_helper_generic.hpp"
 
 #include <cufft.h>
 
@@ -74,7 +75,7 @@ ONEMKL_EXPORT void compute_forward(descriptor_type &desc,
         auto inout_acc = inout.template get_access<sycl::access::mode::read_write>(cgh);
         commit->add_buffer_workspace_dependency_if_rqd("compute_forward", cgh);
 
-        detail::cufft_enqueue_task(cgh, [=](sycl::interop_handle ih) {
+        dft::detail::fft_enqueue_task(cgh, [=](sycl::interop_handle ih) {
             auto stream = detail::setup_stream(func_name, ih, plan);
 
             auto inout_native = reinterpret_cast<fwd<descriptor_type> *>(
@@ -119,7 +120,7 @@ ONEMKL_EXPORT void compute_forward(descriptor_type &desc, sycl::buffer<fwd<descr
         auto out_acc = out.template get_access<sycl::access::mode::read_write>(cgh);
         commit->add_buffer_workspace_dependency_if_rqd("compute_forward", cgh);
 
-        detail::cufft_enqueue_task(cgh, [=](sycl::interop_handle ih) {
+        dft::detail::fft_enqueue_task(cgh, [=](sycl::interop_handle ih) {
             auto stream = detail::setup_stream(func_name, ih, plan);
 
             auto in_native = reinterpret_cast<void *>(
@@ -173,7 +174,7 @@ ONEMKL_EXPORT sycl::event compute_forward(descriptor_type &desc, fwd<descriptor_
         cgh.depends_on(dependencies);
         commit->depend_on_last_usm_workspace_event_if_rqd(cgh);
 
-        detail::cufft_enqueue_task(cgh, [=](sycl::interop_handle ih) {
+        dft::detail::fft_enqueue_task(cgh, [=](sycl::interop_handle ih) {
             auto stream = detail::setup_stream(func_name, ih, plan);
 
             detail::cufft_execute<detail::Direction::Forward, fwd<descriptor_type>>(
@@ -219,7 +220,7 @@ ONEMKL_EXPORT sycl::event compute_forward(descriptor_type &desc, fwd<descriptor_
         cgh.depends_on(dependencies);
         commit->depend_on_last_usm_workspace_event_if_rqd(cgh);
 
-        detail::cufft_enqueue_task(cgh, [=](sycl::interop_handle ih) {
+        dft::detail::fft_enqueue_task(cgh, [=](sycl::interop_handle ih) {
             auto stream = detail::setup_stream(func_name, ih, plan);
 
             detail::cufft_execute<detail::Direction::Forward, fwd<descriptor_type>>(
diff --git a/src/dft/backends/rocfft/backward.cpp b/src/dft/backends/rocfft/backward.cpp
index bdb1c9638..e76437ee2 100644
--- a/src/dft/backends/rocfft/backward.cpp
+++ b/src/dft/backends/rocfft/backward.cpp
@@ -29,6 +29,7 @@
 #include "oneapi/mkl/dft/descriptor.hpp"
 
 #include "execute_helper.hpp"
+#include "../../execute_helper_generic.hpp"
 #include "rocfft_handle.hpp"
 
 #include <rocfft.h>
@@ -78,7 +79,7 @@ ONEMKL_EXPORT void compute_backward(descriptor_type &desc,
         auto inout_acc = inout.template get_access<sycl::access::mode::read_write>(cgh);
         commit->add_buffer_workspace_dependency_if_rqd("compute_backward", cgh);
 
-        detail::rocfft_enqueue_task(cgh, [=](sycl::interop_handle ih) {
+        dft::detail::fft_enqueue_task(cgh, [=](sycl::interop_handle ih) {
             auto stream = detail::setup_stream(func_name, ih, info);
 
             auto inout_native = reinterpret_cast<void *>(
@@ -112,7 +113,7 @@ ONEMKL_EXPORT void compute_backward(descriptor_type &desc,
         auto inout_im_acc = inout_im.template get_access<sycl::access::mode::read_write>(cgh);
         commit->add_buffer_workspace_dependency_if_rqd("compute_backward", cgh);
 
-        detail::rocfft_enqueue_task(cgh, [=](sycl::interop_handle ih) {
+        dft::detail::fft_enqueue_task(cgh, [=](sycl::interop_handle ih) {
             auto stream = detail::setup_stream(func_name, ih, info);
 
             std::array<void *, 2> inout_native{
@@ -146,7 +147,7 @@ ONEMKL_EXPORT void compute_backward(descriptor_type &desc,
         auto out_acc = out.template get_access<sycl::access::mode::read_write>(cgh);
         commit->add_buffer_workspace_dependency_if_rqd("compute_backward", cgh);
 
-        detail::rocfft_enqueue_task(cgh, [=](sycl::interop_handle ih) {
+        dft::detail::fft_enqueue_task(cgh, [=](sycl::interop_handle ih) {
             const std::string func_name = "compute_backward(desc, in, out)";
             auto stream = detail::setup_stream(func_name, ih, info);
 
@@ -181,7 +182,7 @@ ONEMKL_EXPORT void compute_backward(descriptor_type &desc,
         auto out_im_acc = out_im.template get_access<sycl::access::mode::read_write>(cgh);
         commit->add_buffer_workspace_dependency_if_rqd("compute_backward", cgh);
 
-        detail::rocfft_enqueue_task(cgh, [=](sycl::interop_handle ih) {
+        dft::detail::fft_enqueue_task(cgh, [=](sycl::interop_handle ih) {
             const std::string func_name = "compute_backward(desc, in_re, in_im, out_re, out_im)";
             auto stream = detail::setup_stream(func_name, ih, info);
 
@@ -235,7 +236,7 @@ ONEMKL_EXPORT sycl::event compute_backward(descriptor_type &desc, fwd<descriptor
         cgh.depends_on(deps);
         commit->depend_on_last_usm_workspace_event_if_rqd(cgh);
 
-        detail::rocfft_enqueue_task(cgh, [=](sycl::interop_handle ih) {
+        dft::detail::fft_enqueue_task(cgh, [=](sycl::interop_handle ih) {
             auto stream = detail::setup_stream(func_name, ih, info);
 
             void *inout_ptr = inout;
@@ -268,7 +269,7 @@ ONEMKL_EXPORT sycl::event compute_backward(descriptor_type &desc, scalar<descrip
         cgh.depends_on(deps);
         commit->depend_on_last_usm_workspace_event_if_rqd(cgh);
 
-        detail::rocfft_enqueue_task(cgh, [=](sycl::interop_handle ih) {
+        dft::detail::fft_enqueue_task(cgh, [=](sycl::interop_handle ih) {
             auto stream = detail::setup_stream(func_name, ih, info);
 
             std::array<void *, 2> inout_native{ inout_re + offsets[0], inout_im + offsets[0] };
@@ -300,7 +301,7 @@ ONEMKL_EXPORT sycl::event compute_backward(descriptor_type &desc, bwd<descriptor
         cgh.depends_on(deps);
         commit->depend_on_last_usm_workspace_event_if_rqd(cgh);
 
-        detail::rocfft_enqueue_task(cgh, [=](sycl::interop_handle ih) {
+        dft::detail::fft_enqueue_task(cgh, [=](sycl::interop_handle ih) {
             const std::string func_name = "compute_backward(desc, in, out, deps)";
             auto stream = detail::setup_stream(func_name, ih, info);
 
@@ -330,7 +331,7 @@ ONEMKL_EXPORT sycl::event compute_backward(descriptor_type &desc, scalar<descrip
         cgh.depends_on(deps);
         commit->depend_on_last_usm_workspace_event_if_rqd(cgh);
 
-        detail::rocfft_enqueue_task(cgh, [=](sycl::interop_handle ih) {
+        dft::detail::fft_enqueue_task(cgh, [=](sycl::interop_handle ih) {
             const std::string func_name =
                 "compute_backward(desc, in_re, in_im, out_re, out_im, deps)";
             auto stream = detail::setup_stream(func_name, ih, info);
diff --git a/src/dft/backends/rocfft/execute_helper.hpp b/src/dft/backends/rocfft/execute_helper.hpp
index 49c499637..a182546b5 100644
--- a/src/dft/backends/rocfft/execute_helper.hpp
+++ b/src/dft/backends/rocfft/execute_helper.hpp
@@ -98,26 +98,6 @@ inline void execute_checked(const std::string &func, hipStream_t stream, const r
 #endif
 }
 
-/** Wrap interop API to launch interop host task.
- * 
- * @tparam HandlerT The command group handler type
- * @tparam FnT The body of the enqueued task
- *
- * Either uses host task interop API, or enqueue native command extension.
- * This extension avoids host synchronization after 
- * the CUDA call is complete.
- */
-template <typename HandlerT, typename FnT>
-static inline void rocfft_enqueue_task(HandlerT&& cgh, FnT&& f) {
-#ifdef SYCL_EXT_ONEAPI_ENQUEUE_NATIVE_COMMAND
-    cgh.ext_codeplay_enqueue_native_command([=](sycl::interop_handle ih){
-#else
-    cgh.host_task([=](sycl::interop_handle ih){
-#endif
-        f(std::move(ih));
-    });
-}
-
 } // namespace oneapi::mkl::dft::rocfft::detail
 
 #endif
diff --git a/src/dft/backends/rocfft/forward.cpp b/src/dft/backends/rocfft/forward.cpp
index daacc685d..d9a576720 100644
--- a/src/dft/backends/rocfft/forward.cpp
+++ b/src/dft/backends/rocfft/forward.cpp
@@ -30,6 +30,7 @@
 #include "oneapi/mkl/dft/descriptor.hpp"
 
 #include "execute_helper.hpp"
+#include "../../execute_helper_generic.hpp"
 #include "rocfft_handle.hpp"
 
 #include <rocfft.h>
@@ -81,7 +82,7 @@ ONEMKL_EXPORT void compute_forward(descriptor_type &desc,
         auto inout_acc = inout.template get_access<sycl::access::mode::read_write>(cgh);
         commit->add_buffer_workspace_dependency_if_rqd("compute_forward", cgh);
 
-        detail::rocfft_enqueue_task(cgh, [=](sycl::interop_handle ih) {
+        dft::detail::fft_enqueue_task(cgh, [=](sycl::interop_handle ih) {
             auto stream = detail::setup_stream(func_name, ih, info);
 
             auto inout_native = reinterpret_cast<void *>(
@@ -115,7 +116,7 @@ ONEMKL_EXPORT void compute_forward(descriptor_type &desc,
         auto inout_im_acc = inout_im.template get_access<sycl::access::mode::read_write>(cgh);
         commit->add_buffer_workspace_dependency_if_rqd("compute_forward", cgh);
 
-        detail::rocfft_enqueue_task(cgh, [=](sycl::interop_handle ih) {
+        dft::detail::fft_enqueue_task(cgh, [=](sycl::interop_handle ih) {
             auto stream = detail::setup_stream(func_name, ih, info);
 
             std::array<void *, 2> inout_native{
@@ -148,7 +149,7 @@ ONEMKL_EXPORT void compute_forward(descriptor_type &desc, sycl::buffer<fwd<descr
         auto out_acc = out.template get_access<sycl::access::mode::read_write>(cgh);
         commit->add_buffer_workspace_dependency_if_rqd("compute_forward", cgh);
 
-        detail::rocfft_enqueue_task(cgh, [=](sycl::interop_handle ih) {
+        dft::detail::fft_enqueue_task(cgh, [=](sycl::interop_handle ih) {
             const std::string func_name = "compute_forward(desc, in, out)";
             auto stream = detail::setup_stream(func_name, ih, info);
 
@@ -183,7 +184,7 @@ ONEMKL_EXPORT void compute_forward(descriptor_type &desc,
         auto out_im_acc = out_im.template get_access<sycl::access::mode::read_write>(cgh);
         commit->add_buffer_workspace_dependency_if_rqd("compute_forward", cgh);
 
-        detail::rocfft_enqueue_task(cgh, [=](sycl::interop_handle ih) {
+        dft::detail::fft_enqueue_task(cgh, [=](sycl::interop_handle ih) {
             const std::string func_name = "compute_forward(desc, in_re, in_im, out_re, out_im)";
             auto stream = detail::setup_stream(func_name, ih, info);
 
@@ -237,7 +238,7 @@ ONEMKL_EXPORT sycl::event compute_forward(descriptor_type &desc, fwd<descriptor_
         cgh.depends_on(deps);
         commit->depend_on_last_usm_workspace_event_if_rqd(cgh);
 
-        detail::rocfft_enqueue_task(cgh, [=](sycl::interop_handle ih) {
+        dft::detail::fft_enqueue_task(cgh, [=](sycl::interop_handle ih) {
             auto stream = detail::setup_stream(func_name, ih, info);
 
             void *inout_ptr = inout;
@@ -269,7 +270,7 @@ ONEMKL_EXPORT sycl::event compute_forward(descriptor_type &desc, scalar<descript
     sycl::event sycl_event = queue.submit([&](sycl::handler &cgh) {
         cgh.depends_on(deps);
         commit->depend_on_last_usm_workspace_event_if_rqd(cgh);
-        detail::rocfft_enqueue_task(cgh, [=](sycl::interop_handle ih) {
+        dft::detail::fft_enqueue_task(cgh, [=](sycl::interop_handle ih) {
             auto stream = detail::setup_stream(func_name, ih, info);
 
             std::array<void *, 2> inout_native{ inout_re + offsets[0], inout_im + offsets[0] };
@@ -300,7 +301,7 @@ ONEMKL_EXPORT sycl::event compute_forward(descriptor_type &desc, fwd<descriptor_
         cgh.depends_on(deps);
         commit->depend_on_last_usm_workspace_event_if_rqd(cgh);
 
-        detail::rocfft_enqueue_task(cgh, [=](sycl::interop_handle ih) {
+        dft::detail::fft_enqueue_task(cgh, [=](sycl::interop_handle ih) {
             const std::string func_name = "compute_forward(desc, in, out, deps)";
             auto stream = detail::setup_stream(func_name, ih, info);
 
@@ -330,7 +331,7 @@ ONEMKL_EXPORT sycl::event compute_forward(descriptor_type &desc, scalar<descript
         cgh.depends_on(deps);
         commit->depend_on_last_usm_workspace_event_if_rqd(cgh);
 
-        detail::rocfft_enqueue_task(cgh, [=](sycl::interop_handle ih) {
+        dft::detail::fft_enqueue_task(cgh, [=](sycl::interop_handle ih) {
             const std::string func_name =
                 "compute_forward(desc, in_re, in_im, out_re, out_im, deps)";
             auto stream = detail::setup_stream(func_name, ih, info);
diff --git a/src/dft/execute_helper_generic.hpp b/src/dft/execute_helper_generic.hpp
new file mode 100644
index 000000000..519f6fda6
--- /dev/null
+++ b/src/dft/execute_helper_generic.hpp
@@ -0,0 +1,53 @@
+/*******************************************************************************
+* Copyright Codeplay Software Ltd.
+*
+* Licensed under the Apache License, Version 2.0 (the "License");
+* you may not use this file except in compliance with the License.
+* You may obtain a copy of the License at
+*
+* http://www.apache.org/licenses/LICENSE-2.0
+*
+* Unless required by applicable law or agreed to in writing,
+* software distributed under the License is distributed on an "AS IS" BASIS,
+* WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
+* See the License for the specific language governing permissions
+* and limitations under the License.
+*
+*
+* SPDX-License-Identifier: Apache-2.0
+*******************************************************************************/
+
+#ifndef _ONEMKL_DFT_SRC_CUFFT_EXECUTE_GENERIC_HPP_
+#define _ONEMKL_DFT_SRC_CUFFT_EXECUTE_GENERIC_HPP_
+
+#if __has_include(<sycl/sycl.hpp>)
+#include <sycl/sycl.hpp>
+#else
+#include <CL/sycl.hpp>
+#endif
+
+namespace oneapi::mkl::dft::detail {
+
+/** Wrap interop API to launch interop host task.
+ * 
+ * @tparam HandlerT The command group handler type
+ * @tparam FnT The body of the enqueued task
+ *
+ * Either uses host task interop API, or enqueue native command extension.
+ * This extension avoids host synchronization after 
+ * the native call is complete.
+ */
+template <typename HandlerT, typename FnT>
+static inline void fft_enqueue_task(HandlerT&& cgh, FnT&& f) {
+#ifdef SYCL_EXT_ONEAPI_ENQUEUE_NATIVE_COMMAND
+    cgh.ext_codeplay_enqueue_native_command([=](sycl::interop_handle ih){
+#else
+    cgh.host_task([=](sycl::interop_handle ih){
+#endif
+        f(std::move(ih));
+    });
+}
+
+} // namespace oneapi::mkl::dft::detail
+
+#endif

From 210dade6a91dbc22cd2b7b4348b1e9041191b481 Mon Sep 17 00:00:00 2001
From: JackAKirk <jack.kirk@codeplay.com>
Date: Fri, 11 Oct 2024 07:59:20 -0700
Subject: [PATCH 09/10] Use more sensible header macro name

Signed-off-by: JackAKirk <jack.kirk@codeplay.com>
---
 src/dft/backends/cufft/execute_helper.hpp  | 2 +-
 src/dft/backends/rocfft/execute_helper.hpp | 2 +-
 src/dft/execute_helper_generic.hpp         | 6 +++---
 3 files changed, 5 insertions(+), 5 deletions(-)

diff --git a/src/dft/backends/cufft/execute_helper.hpp b/src/dft/backends/cufft/execute_helper.hpp
index bbe32c146..11f32c71b 100644
--- a/src/dft/backends/cufft/execute_helper.hpp
+++ b/src/dft/backends/cufft/execute_helper.hpp
@@ -149,4 +149,4 @@ inline CUstream setup_stream(const std::string &func, sycl::interop_handle ih, c
 
 } // namespace oneapi::mkl::dft::cufft::detail
 
-#endif
+#endif // _ONEMKL_DFT_SRC_CUFFT_EXECUTE_HPP_
diff --git a/src/dft/backends/rocfft/execute_helper.hpp b/src/dft/backends/rocfft/execute_helper.hpp
index a182546b5..51af1f6be 100644
--- a/src/dft/backends/rocfft/execute_helper.hpp
+++ b/src/dft/backends/rocfft/execute_helper.hpp
@@ -100,4 +100,4 @@ inline void execute_checked(const std::string &func, hipStream_t stream, const r
 
 } // namespace oneapi::mkl::dft::rocfft::detail
 
-#endif
+#endif // _ONEMKL_DFT_SRC_ROCFFT_EXECUTE_HELPER_HPP_
diff --git a/src/dft/execute_helper_generic.hpp b/src/dft/execute_helper_generic.hpp
index 519f6fda6..4a1747a42 100644
--- a/src/dft/execute_helper_generic.hpp
+++ b/src/dft/execute_helper_generic.hpp
@@ -17,8 +17,8 @@
 * SPDX-License-Identifier: Apache-2.0
 *******************************************************************************/
 
-#ifndef _ONEMKL_DFT_SRC_CUFFT_EXECUTE_GENERIC_HPP_
-#define _ONEMKL_DFT_SRC_CUFFT_EXECUTE_GENERIC_HPP_
+#ifndef _ONEMKL_DFT_SRC_EXECUTE_GENERIC_HPP_
+#define _ONEMKL_DFT_SRC_EXECUTE_GENERIC_HPP_
 
 #if __has_include(<sycl/sycl.hpp>)
 #include <sycl/sycl.hpp>
@@ -50,4 +50,4 @@ static inline void fft_enqueue_task(HandlerT&& cgh, FnT&& f) {
 
 } // namespace oneapi::mkl::dft::detail
 
-#endif
+#endif // _ONEMKL_DFT_SRC_EXECUTE_GENERIC_HPP_

From d76a67b268bd8fb50e30ceb87d65578391ef47d6 Mon Sep 17 00:00:00 2001
From: JackAKirk <jack.kirk@codeplay.com>
Date: Fri, 11 Oct 2024 08:08:11 -0700
Subject: [PATCH 10/10] Make header macro naming consistent

Signed-off-by: JackAKirk <jack.kirk@codeplay.com>
---
 src/dft/backends/cufft/execute_helper.hpp  | 6 +++---
 src/dft/backends/rocfft/execute_helper.hpp | 6 +++---
 src/dft/execute_helper_generic.hpp         | 6 +++---
 3 files changed, 9 insertions(+), 9 deletions(-)

diff --git a/src/dft/backends/cufft/execute_helper.hpp b/src/dft/backends/cufft/execute_helper.hpp
index 11f32c71b..7b7d946db 100644
--- a/src/dft/backends/cufft/execute_helper.hpp
+++ b/src/dft/backends/cufft/execute_helper.hpp
@@ -17,8 +17,8 @@
 * SPDX-License-Identifier: Apache-2.0
 *******************************************************************************/
 
-#ifndef _ONEMKL_DFT_SRC_CUFFT_EXECUTE_HPP_
-#define _ONEMKL_DFT_SRC_CUFFT_EXECUTE_HPP_
+#ifndef _ONEMKL_DFT_SRC_EXECUTE_HELPER_CUFFT_HPP_
+#define _ONEMKL_DFT_SRC_EXECUTE_HELPER_CUFFT_HPP_
 
 #if __has_include(<sycl/sycl.hpp>)
 #include <sycl/sycl.hpp>
@@ -149,4 +149,4 @@ inline CUstream setup_stream(const std::string &func, sycl::interop_handle ih, c
 
 } // namespace oneapi::mkl::dft::cufft::detail
 
-#endif // _ONEMKL_DFT_SRC_CUFFT_EXECUTE_HPP_
+#endif // _ONEMKL_DFT_SRC_EXECUTE_HELPER_CUFFT_HPP_
diff --git a/src/dft/backends/rocfft/execute_helper.hpp b/src/dft/backends/rocfft/execute_helper.hpp
index 51af1f6be..78663a090 100644
--- a/src/dft/backends/rocfft/execute_helper.hpp
+++ b/src/dft/backends/rocfft/execute_helper.hpp
@@ -17,8 +17,8 @@
 * SPDX-License-Identifier: Apache-2.0
 *******************************************************************************/
 
-#ifndef _ONEMKL_DFT_SRC_ROCFFT_EXECUTE_HELPER_HPP_
-#define _ONEMKL_DFT_SRC_ROCFFT_EXECUTE_HELPER_HPP_
+#ifndef _ONEMKL_DFT_SRC_EXECUTE_HELPER_ROCFFT_HPP_
+#define _ONEMKL_DFT_SRC_EXECUTE_HELPER_ROCFFT_HPP_
 
 #if __has_include(<sycl/sycl.hpp>)
 #include <sycl/sycl.hpp>
@@ -100,4 +100,4 @@ inline void execute_checked(const std::string &func, hipStream_t stream, const r
 
 } // namespace oneapi::mkl::dft::rocfft::detail
 
-#endif // _ONEMKL_DFT_SRC_ROCFFT_EXECUTE_HELPER_HPP_
+#endif // _ONEMKL_DFT_SRC_EXECUTE_HELPER_ROCFFT_HPP_
diff --git a/src/dft/execute_helper_generic.hpp b/src/dft/execute_helper_generic.hpp
index 4a1747a42..22fe0cb33 100644
--- a/src/dft/execute_helper_generic.hpp
+++ b/src/dft/execute_helper_generic.hpp
@@ -17,8 +17,8 @@
 * SPDX-License-Identifier: Apache-2.0
 *******************************************************************************/
 
-#ifndef _ONEMKL_DFT_SRC_EXECUTE_GENERIC_HPP_
-#define _ONEMKL_DFT_SRC_EXECUTE_GENERIC_HPP_
+#ifndef _ONEMKL_DFT_SRC_EXECUTE_HELPER_GENERIC_HPP_
+#define _ONEMKL_DFT_SRC_EXECUTE_HELPER_GENERIC_HPP_
 
 #if __has_include(<sycl/sycl.hpp>)
 #include <sycl/sycl.hpp>
@@ -50,4 +50,4 @@ static inline void fft_enqueue_task(HandlerT&& cgh, FnT&& f) {
 
 } // namespace oneapi::mkl::dft::detail
 
-#endif // _ONEMKL_DFT_SRC_EXECUTE_GENERIC_HPP_
+#endif // _ONEMKL_DFT_SRC_EXECUTE_HELPER_GENERIC_HPP_