sycl: bound in-flight expert matmuls in mul_mat_id (fix MoE OUT_OF_RESOURCES on Intel iGPU) - #24635
sycl: bound in-flight expert matmuls in mul_mat_id (fix MoE OUT_OF_RESOURCES on Intel iGPU)#24635mayerwin wants to merge 1 commit into
Conversation
github-actions
Bot
added
ggml
|
Hi @mayerwin, thanks for your contribution! Per our contribution guidelines, the automated PR checker found the following issue(s) that need your attention:
Please note that maintainers reserve the right to make final decisions on PRs. If you believe there is a mistake, please comment below. |
|
@mayerwin I'm interesting in this issue.
I didn't reproduce this issue by Level Zero on B60 on Ubuntu 24.04.
Thank you! |
…URCES 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.
c413c1e to
4389502
Compare
|
Thanks @arthw, this is really helpful. You're right on several points and I've reworked the patch. Single
Both aborted with
if (stream->get_backend() != sycl::backend::ext_oneapi_level_zero) {
SYCL_CHECK(CHECK_TRY_ERROR(stream->wait()));
}Level Zero (points 1 and 2): I'm on OpenCL because Repro: Arc 140T iGPU (Arrow Lake-P / Xe2-LPG), Ubuntu 24.04, compute-runtime 24.39.31294; any MoE with experts on GPU: Perf table (point 3): you're right, that was apples to oranges (experts-CPU vs experts-GPU). Removed it. This isn't a perf optimization; it makes a config that currently hard-crashes (full-GPU MoE on OpenCL) run at all, so there's no same-config without-flush number. One note: a reviewer flagged that a single post-loop drain (vs an in-loop chunk) could in theory still overflow for one op with far more experts than tested, on a weaker OpenCL device. The 512-expert pass is reassuring, but I'm glad to add a small in-loop chunk (still gated) if you'd prefer the margin. |
|
@mayerwin
That means you still can run on Level-Zero even if Level zero API for memory malloc is not supported. And, I find the UT cases of mul_mat_id is unstable recently on dGPU. So the PR will be pending to review. Thank you! |
|
Thanks, that matches what I found, and good to confirm the Level Zero runtime is still used regardless of the malloc path. Makes sense to wait for your mul_mat_id investigation. One possibly-related data point: while validating this on the Arc 140T (OpenCL) I hit an intermittent failure in Happy to wait, and glad to re-test this PR and your mul_mat_id fix on the iGPU once it lands. |
|
@mayerwin You could test on Level-Zero running time. Thank you! |
|
@mayerwin Could you try your test by this cmd? Thank you! |
|
Update after more testing on the Arc 140T iGPU (OpenCL), and I'll close this. On On the original Since I cannot demonstrate the crash or that my change is needed, and it adds a host sync that is pure overhead when the issue is absent, I am closing this. Thanks for the review and the pointers. I will reopen with a solid repro if I hit it again, and I'm glad to test your |
3 participants
Summary
On the SYCL OpenCL backend (Intel Arc iGPU), running MoE expert tensors on the GPU (
-ngl 99, no-ot exps=CPU) aborts during multi-token prompt processing:I hit this on every MoE model I tried with experts on GPU: gemma-4-26b-a4b, qwen3.6-35b-a3b, qwen3-coder-next-80b. This patch makes that configuration run; it is a correctness/availability fix, not a performance optimization.
Root cause
ggml_sycl_mul_mat_idenqueues one matmul per expert (n_asup to hundreds). On a single deviceggml_sycl_op_mul_matdoes no host wait, and the node-submit loop never waits between graph ops, so the per-expert matmuls accumulate across ops until the nextmul_mat_id's start-of-function wait, exceeding the OpenCL UR adapter's in-flight command/event budget.It is not a single bad kernel: a 128-expert
mul_mat_idin isolation (test-backend-ops) does not crash; it only manifests across the full model graph. Anything that throttles submission avoids it (-ub 1, or even heavy debug logging).Fix
Drain the queue once after the expert loop, gated to the non-Level-Zero backend (no-op on Level Zero):
This caps in-flight expert matmuls at a single op's worth. It is numerically inert (it only synchronizes).
Validation (Intel Arc 140T iGPU, OpenCL UR)
No crash, on configs that previously aborted within seconds:
test-backend-ops -o MUL_MAT_IDpasses. Since the fix only adds a host wait, it cannot change results.Repro
Disclosure: I used an AI coding assistant to help with this. The code change is small; the bulk of the work was diagnosis (reproducing the abort, ruling out the work-group / single-kernel theories, and tracing it to async submission pile-up). The investigation and on-hardware validation are mine.