From 4389502f0cdcd7172c9882b9e09f03cc640a6726 Mon Sep 17 00:00:00 2001 From: Erwin Mayer Date: Sun, 14 Jun 2026 19:19:15 -1000 Subject: [PATCH] sycl: drain queue after mul_mat_id expert loop to fix MoE OUT_OF_RESOURCES on Intel iGPU Running MoE expert tensors on the GPU aborts during multi-token prompt processing on the SYCL OpenCL backend (Intel Arc iGPU) with UR_RESULT_ERROR_OUT_OF_RESOURCES, surfacing at the stream->wait() in ggml_sycl_mul_mat_id. ggml_sycl_mul_mat_id enqueues one matmul per expert (n_as up to hundreds). ggml_sycl_op_mul_mat does no host wait on a single device and the node-submit loop never waits between ops, so without a drain the per-expert matmuls accumulate across ops until the next mul_mat_id's wait, exceeding the OpenCL UR adapter's in-flight budget. A single 128-expert mul_mat_id in isolation does not crash; it only manifests across the full model graph. Drain the queue once after the expert loop, gated to the non-Level-Zero backend (no-op on Level Zero, the default Intel path). This caps in-flight expert matmuls at a single op's worth. Validated with no crash on Arc 140T (OpenCL) for gemma-4-26b-a4b (128 experts) and qwen3-coder-next-80b (512 experts); full-GPU MoE now runs where it previously aborted. --- ggml/src/ggml-sycl/ggml-sycl.cpp | 10 ++++++++++ 1 file changed, 10 insertions(+) diff --git a/ggml/src/ggml-sycl/ggml-sycl.cpp b/ggml/src/ggml-sycl/ggml-sycl.cpp index 3f246e8672d..ee0593780f0 100644 --- a/ggml/src/ggml-sycl/ggml-sycl.cpp +++ b/ggml/src/ggml-sycl/ggml-sycl.cpp @@ -4393,6 +4393,16 @@ static void ggml_sycl_mul_mat_id(ggml_backend_sycl_context & ctx, ggml_sycl_mul_mat(ctx, &src0_row, &src1_row, &dst_row); } + // The OpenCL UR adapter aborts with UR_RESULT_ERROR_OUT_OF_RESOURCES on this MoE op: it + // enqueues one matmul per expert (n_as up to hundreds), ggml_sycl_op_mul_mat does no host + // wait on a single device, and the node-submit loop never waits between ops, so without a + // drain here expert matmuls accumulate across ops until the next mul_mat_id's wait, + // exceeding the adapter's in-flight budget. One drain after the loop bounds this to a + // single mul_mat_id's experts. Level Zero did not hit this in our testing, so skip it there. + if (stream->get_backend() != sycl::backend::ext_oneapi_level_zero) { + SYCL_CHECK(CHECK_TRY_ERROR(stream->wait())); + } + { sycl::range<3> block_dims(1, 1, std::min((unsigned int)ne0, max_work_group_size)); sycl::range<3> grid_dims(1, 1, n_routed_rows);