Feature Request: SYCL per-instance queue, device-side events, async copy

[mobilinkd]

Prerequisites

• I am running the latest code. Mention the version if possible as well.
• I carefully followed the README.md.
• I searched using keywords relevant to my issue to make sure that I am creating a new issue that is not already open (or closed).
• I reviewed the Discussions, and have a new and useful enhancement to share.

Feature Description

The SYCL backend aliases the dpct default queue for every context, so a
second backend instance (e.g. the scheduler's copy backend) shares the
compute queue and no copy/compute overlap is possible.

Allow each context to own its queue (q_owned, in-order); other devices in
tensor-split mode still use the default queue.

event_wait waits device-side via ext_oneapi_submit_barrier instead of
host-blocking sycl_event->wait(), mirroring cudaStreamWaitEvent and letting
the compute stream defer to a copy event without serializing the pipeline.

Motivation

Without this feature, features such as #21067 and others which attempt to overlap expert tensor prefetch with compute fail to overlap on SYCL (measured on Intel Arc Pro B70 with vtune). No speedup is achieved. In fact, performance regresses slightly because these features blindly stream all experts, some of which are not needed.

The numbers (DSv4F IQ2_XXS, pp2048, branch prefetch-08-13-test off master a94d563, -ngl 44 --n-cpu-moe 43 -lm none -ub
2048, GGML_SCHED_PREFETCH_EXPERTS=4, GPU idle-gated):

     stack                                     pp2048    vs baseline                                                                     
     master + 2 GiB cap (baseline)             167.32        --                                                                          
     + dpct glue, NO prefetch (8027c1c45)      167.50      +0.1%   neutral                                                               
     + prefetch scheduler, NO glue             160.27      -4.2%   REGRESSION                                                            
     + prefetch + glue                         193.38     +15.6%                                                                         

NOTE: #21067 appears to be CUDA-specific and does not improve performance on SYCL even with these patches. A re-worked patch which does the same thing and which is backend agnostic does speed up the performance on SYCL. The above numbers are based on a custom adaptation of work originally done by @thecodacus (see #25859)

Possible Implementation

@arthw

The SYCL backend aliased the dpct default queue for every context, so a
second backend instance (e.g. the scheduler's copy backend) shared the
compute queue and no copy/compute overlap was possible. Each context now
owns its queue (q_owned, in-order); other devices in tensor-split mode
still use the default queue.

event_wait now waits device-side via ext_oneapi_submit_barrier instead of
host-blocking sycl_event->wait(), mirroring cudaStreamWaitEvent and letting
the compute stream defer to a copy event without serializing the pipeline.

cpy_tensor_async wired with the 4-arg backend interface (was NULL):
same-backend = queue memcpy, cross-backend = memcpy on src stream +
barrier on dst stream.

Assisted-by: Hermes Agent

 ggml/src/ggml-sycl/common.hpp    | 17 +++++++-
 ggml/src/ggml-sycl/ggml-sycl.cpp | 67 ++++++++++++++++++++------------
 2 files changed, 58 insertions(+), 26 deletions(-)

diff --git a/ggml/src/ggml-sycl/common.hpp b/ggml/src/ggml-sycl/common.hpp
index 355dd442b..86517bf3d 100644
--- a/ggml/src/ggml-sycl/common.hpp
+++ b/ggml/src/ggml-sycl/common.hpp
@@ -340,15 +340,30 @@ struct ggml_backend_sycl_context {
 
     queue_ptr qptrs[GGML_SYCL_MAX_DEVICES][GGML_SYCL_MAX_STREAMS] = { { nullptr } };
 
+    // queue owned by this backend instance. a second instance (e.g. the copy
+    // backend created by the scheduler prefetch path) gets its own queue, which
+    // is what makes copy/compute overlap possible. mirrors the per-instance
+    // stream of the CUDA backend. dpct registers it, so queues_wait_and_throw
+    // covers it during teardown.
+    sycl::queue q_owned;
+
     explicit ggml_backend_sycl_context(int device) :
         device(device),
         name(GGML_SYCL_NAME + std::to_string(device)) {
         opt_feature = ggml_sycl_info().devices[device].opt_feature;
+        q_owned = dpct::get_device(device).create_queue();
+        qptrs[device][0] = &q_owned;
     }
 
     queue_ptr stream(int device, int stream) {
         if (qptrs[device][stream] == nullptr) {
-            qptrs[device][stream] = &(dpct::get_device(device).default_queue());
+            if (device == this->device) {
+                // this backend's own queue, created in the constructor
+                qptrs[device][stream] = &q_owned;
+            } else {
+                // other devices (tensor-split mode) share the dpct default queue
+                qptrs[device][stream] = &(dpct::get_device(device).default_queue());
+            }
         }
         return qptrs[device][stream];
     }
diff --git a/ggml/src/ggml-sycl/ggml-sycl.cpp b/ggml/src/ggml-sycl/ggml-sycl.cpp
index 4091f73a4..8a05e1799 100644
--- a/ggml/src/ggml-sycl/ggml-sycl.cpp
+++ b/ggml/src/ggml-sycl/ggml-sycl.cpp
@@ -5798,29 +5798,43 @@ catch (sycl::exception const &exc) {
   std::exit(1);
 }
 
-static bool ggml_backend_sycl_cpy_tensor_async(ggml_backend_t backend,
-                                               const ggml_tensor *src,
-                                               ggml_tensor *dst) try {
-    ggml_backend_sycl_context * sycl_ctx = (ggml_backend_sycl_context *)backend->context;
-    bool is_cpy_supported                = dst->buffer->buft == ggml_backend_sycl_buffer_type(sycl_ctx->device) &&
-                            ggml_backend_buffer_is_sycl(src->buffer);
-    GGML_SYCL_DEBUG("[SYCL] call %s", __func__);
-    GGML_SYCL_DEBUG("%s", debug_get_tensor_str(": dst", dst).c_str());
-    GGML_SYCL_DEBUG("%s", debug_get_tensor_str(" src", src).c_str());
-    GGML_SYCL_DEBUG(" is_cpy_supported=%d\n", is_cpy_supported);
-    if (is_cpy_supported) {
-        /*
-        DPCT1009:215: SYCL uses exceptions to report errors and does not use the
-        error codes. The original code was commented out and a warning string
-        was inserted. You need to rewrite this code.
-        */
-        const queue_ptr stream = sycl_ctx->stream(sycl_ctx->device, 0);
-        SYCL_CHECK(CHECK_TRY_ERROR((stream)->memcpy(
-            dst->data, src->data, ggml_nbytes(dst))));
-        return true;
+static bool ggml_backend_sycl_cpy_tensor_async(ggml_backend_t backend_src, ggml_backend_t backend_dst, const ggml_tensor *src, ggml_tensor *dst) try {
+    ggml_backend_buffer_t buf_src = src->view_src ? src->view_src->buffer : src->buffer;
+    ggml_backend_buffer_t buf_dst = dst->view_src ? dst->view_src->buffer : dst->buffer;
+
+    if (!ggml_backend_is_sycl(backend_src) || !ggml_backend_is_sycl(backend_dst)) {
+        return false;
     }
 
-    return false;
+    if (!ggml_backend_buffer_is_sycl(buf_src) || !ggml_backend_buffer_is_sycl(buf_dst)) {
+        return false;
+    }
+
+    ggml_backend_sycl_context * sycl_ctx_src = (ggml_backend_sycl_context *)backend_src->context;
+    ggml_backend_sycl_context * sycl_ctx_dst = (ggml_backend_sycl_context *)backend_dst->context;
+
+    ggml_backend_sycl_buffer_context * buf_ctx_src = (ggml_backend_sycl_buffer_context *) buf_src->context;
+    ggml_backend_sycl_buffer_context * buf_ctx_dst = (ggml_backend_sycl_buffer_context *) buf_dst->context;
+
+    if (sycl_ctx_src->device != buf_ctx_src->device || sycl_ctx_dst->device != buf_ctx_dst->device) {
+        GGML_LOG_DEBUG("%s: backend and buffer devices do not match\n", __func__);
+        return false;
+    }
+
+    if (backend_src != backend_dst) {
+        // copy on src stream
+        const queue_ptr stream_src = sycl_ctx_src->stream(sycl_ctx_src->device, 0);
+        sycl::event copy_event = stream_src->memcpy(dst->data, src->data, ggml_nbytes(dst));
+
+        // wait on dst stream for the copy to complete
+        const queue_ptr stream_dst = sycl_ctx_dst->stream(sycl_ctx_dst->device, 0);
+        stream_dst->ext_oneapi_submit_barrier({copy_event});
+    } else {
+        // src and dst are on the same backend
+        const queue_ptr stream = sycl_ctx_src->stream(sycl_ctx_src->device, 0);
+        stream->memcpy(dst->data, src->data, ggml_nbytes(dst));
+    }
+    return true;
 }
 catch (sycl::exception const &exc) {
   std::cerr << exc.what() << "Exception caught at file:" << __FILE__
@@ -6120,7 +6134,12 @@ static void ggml_backend_sycl_event_wait(ggml_backend_t backend, ggml_backend_ev
     sycl::event* sycl_event = static_cast<sycl::event*>(event->context);
 
     if (ggml_backend_is_sycl(backend)) {
-        SYCL_CHECK(CHECK_TRY_ERROR(sycl_event->wait()));
+        // device-side wait: make this backend's queue wait for the event without
+        // blocking the host (mirror of cudaStreamWaitEvent). required by the
+        // scheduler prefetch pipeline to overlap copy and compute.
+        ggml_backend_sycl_context * sycl_ctx = (ggml_backend_sycl_context *)backend->context;
+        const queue_ptr stream = sycl_ctx->stream(sycl_ctx->device, 0);
+        SYCL_CHECK(CHECK_TRY_ERROR(stream->ext_oneapi_submit_barrier({*sycl_event})));
     } else
         GGML_ABORT("fatal error");
 } catch (sycl::exception const& exc) {
@@ -6136,9 +6155,7 @@ static ggml_backend_i ggml_backend_sycl_interface = {
     /* .get_tensor_async        = */ ggml_backend_sycl_get_tensor_async,
     /* .set_tensor_2d_async     = */ NULL,
     /* .get_tensor_2d_async     = */ NULL,
-    /* .cpy_tensor_async        = */ NULL, // ggml_backend_sycl_cpy_tensor_async,
-                                           // // TODO: update for the new
-                                           // interface
+    /* .cpy_tensor_async        = */ ggml_backend_sycl_cpy_tensor_async,
     /* .synchronize             = */ ggml_backend_sycl_synchronize,
     /* .graph_plan_create       = */ NULL,
     /* .graph_plan_free         = */ NULL,