Skip to content

feat: support Qwen3.5 eager-mode on NPU with Python executor. - #2264

Open
Wang-1F wants to merge 1 commit into
xLLM-AI:mainfrom
Wang-1F:rebase_main
Open

feat: support Qwen3.5 eager-mode on NPU with Python executor.#2264
Wang-1F wants to merge 1 commit into
xLLM-AI:mainfrom
Wang-1F:rebase_main

Conversation

@Wang-1F

@Wang-1F Wang-1F commented Aug 18, 2026

Copy link
Copy Markdown
Collaborator

Summary

Support Qwen3.5 on NPU via the Python executor path by replacing all PyTorch small-op kernel stubs with high-performance NPU operator calls.

What this PR does

  • Registers 8 new torch.ops.xllm_ops operators for the Python executor's GDN (Gated Delta Net) layer on NPU:

    • rms_norm_gated — fused RMSNorm + SiLU gate (CANN aclnnLayerNormFwd)
    • l2_norm — L2 normalization (Triton-on-Ascend)
    • causal_conv1d_update — decode-path causal convolution (TileLang)
    • causal_conv1d_prefill — prefill-path causal convolution (CANN aclnnCausalConv1d)
    • causal_conv1d_qkv_prefill — fused conv + QKV split + L2norm (CANN aclnnCausalConv1dQkv)
    • chunk_gated_delta_rule — prefill delta-rule recurrence (CANN aclnnMegaChunkGdn)
    • fused_gdn_gating — decay gate computation (TileLang)
    • fused_sigmoid_gating_delta_rule_decode — fused decode recurrence (Triton-on-Ascend)
  • Refactors ops_api.cpp NPU causal_conv1d_update implementation into a reusable tilelang wrapper (causal_conv1d_update_wrapper.cpp), shared by both the C++ model layer and the Python torch.ops dispatch.

  • Introduces gdn_prefill_prepare as a unified interface for both CUDA and NPU prefill paths, encapsulating platform-specific fusion strategies.

  • Enables qwen3_5 in model_platform_support.py for NPU.

Precision

  • Layer 0 hidden_states: bit-exact match with C++ native path
  • Layer 0 core_attn_out: max diff = 2.4e-4 (1 ULP in bf16)
  • End-to-end output tokens: identical to C++ native for short prompts (temperature=0)

Test plan

  • xllm serve --model_impl=python starts successfully with Qwen3.5-4B on NPU
  • End-to-end inference produces correct, coherent output (no garbled text)
  • Output token sequence matches C++ native path for short prompts
  • Multiple prompts (Chinese/English, short/long) verified correct
  • Run python_executor_test after fixing _EXPORTS sync
  • Performance benchmarking vs C++ native path

Copy link
Copy Markdown
Collaborator

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

为什么需要这个?

@yingxudeng

Copy link
Copy Markdown
Collaborator
image 格式检查没过

params.block_idx_last_scheduled_token,
params.initial_state_idx);
#elif defined(USE_NPU)
const bool has_silu = params.activation;

Copy link
Copy Markdown
Collaborator

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

是否会影响libtorch 组图的 qwen3.5

@zhang-minchao zhang-minchao changed the title feat: support Qwen3.5 eagle-mode on NPU with Python executor. feat: support Qwen3.5 eager-mode on NPU with Python executor. Aug 19, 2026
Sign up for free to join this conversation on GitHub. Already have an account? Sign in to comment

Labels

None yet

Projects

None yet

Development

Successfully merging this pull request may close these issues.

3 participants