[wdziurdz]

moe_align_block_size creates a sycl::queue on every call that is never used; the kernels are submitted to the current XPU stream. Creating it costs 1-2 ms of host time per call, and SGLang's Triton MoE runner calls this op once per MoE layer per forward, so eager decode becomes host bound.

DeepSeek on Arc Pro B60, SGLang with --attention-backend triton --moe-runner-backend triton, sglang.benchmark.one_batch 1024 in / 32 out:

model	metric	before	after
DeepSeek-V2-Lite-Chat-FP8	decode step, bs1	95.7 ms	60.3 ms
DeepSeek-V2-Lite-Chat-FP8	decode step, bs16	96.5 ms	65.8 ms
DeepSeek-OCR	decode step, bs16	27.8 ms	13.1 ms

Test

pytest tests/test_moe_align.py

5168 passed.

[Copilot]

🟢 Approval recommended

No active review issues are recorded for the removal of unused per-call SYCL queue construction.

No open issues 🎉

What changed in this PR

Removes unused per-call SYCL queue construction from moe_align_block_size to reduce host overhead in MoE inference. Kernel submissions remain on the current XPU stream.

Changes:

• Deletes the unused sycl::queue variable without changing kernel logic.
• Reports improved DeepSeek decode latency and 5,168 passing MoE alignment tests.

File	Description
src/​sycl/​MoEAlign.cpp	Removes unused queue construction from moe_align_block_size.

_{Review effort: Balanced}

[wdziurdz]

cc @airMeng @sunjiweiswift could you take a look?