feat: 昇腾 NPU 支持 — MossVL 推理（sglang serving + 独立推理脚本）

[Joiin0392]

概述

为 MOSS-VL 增加华为昇腾（Ascend）NPU 推理支持。包含两条路径：

1. sglang serving（高性能，OpenAI 兼容 API）
2. 独立推理脚本（纯 transformers，不依赖 sglang）

所有改动均以 NPU 门控，非 NPU 环境走原有代码路径。

改动内容（15 文件，+310/-22）

新增文件

文件	作用
device_utils.py	CUDA/NPU 双后端设备抽象（设备检测、device 字符串、attn 实现解析、同步原语）——两个平台都是一等公民
requirements-npu.txt	NPU 环境依赖闭环（torch-npu 2.8.0 / transformers 4.57.1 / torchcodec 等），上游 pyproject.toml 零改动（CUDA 依赖原样）

独立推理脚本（上游入口的 NPU 适配，各 ~7 行）

• inference/run_inference.py、realtime_inference/run_online_inference.py：
  • attn_implementation 从硬编码 flash_attention_2（CUDA-only）改为经 device_utils.resolve_attn_impl() 解析（NPU → eager，CUDA → 原值不变）
  • --attention-backend 默认值 flash_attention_2 → auto

sglang 修复（9 文件）

文件	修复
server_args.py	NPU 上 MossVL 默认 ascend attention 后端；flashinfer 断言放宽（device != 'npu'；NPU 上 radix cache 自动禁用（昇腾 paged-attention KV 布局敏感，temperature=0 下解码有 ~20% 非确定性失败）
models/moss_vl.py	非 flashinfer prefill 后端时跳过 FlashInfer 打包 cross-attention mask（ascend 退回标准全可见 cross-attention——正确且非降级）
models/qwen3.py + transcription_adapters/__init__.py	启动阻断修复：sglang 启动时 import 全部模型注册表，上游 qwen3.py 顶层硬 import sgl_kernel_npu.norm.split_qkv_rmsnorm_rope，当环境的 triton-ascend 缺 triton.language.extra.cann.extension 时整个 sglang 启动崩溃（跑 MossVL 也一样炸）。降级为 try/except graceful fallback——Qwen3 缺内核时退回慢路径，其它模型不再受牵连。任何 sgl_kernel_npu 安装不完整的环境都会撞上此问题
ascend_backend.py	解码 cross-attention 处理 k/v 为 None（仅 NPU 硬件后端路径，CUDA 不加载）
profile_utils.py	torch_npu._apply_all_patches() 加 hasattr 守卫（torch-npu 2.10+ API 改名兼容）
schedule_batch.py	丢弃落在视觉 token 区间内的 encoder-prefix 缓存匹配（paged KV 下违反 encoder 整体单元语义；CUDA 默认 page_size=1 时为死代码）
scheduler.py / scheduler_profiler_mixin.py / disaggregation/decode.py / dllm/mixin/scheduler.py	NPU 设备类型守卫的小型适配

使用方式

sglang serving（NPU）

pip install -e sglang/python --no-deps
pip install -r requirements-npu.txt

python -m sglang.launch_server \
  --model-path MOSS-VL-Instruct-0708 \
  --device npu --base-gpu-id 0 \
  --attention-backend ascend \
  --prefill-attention-backend ascend \
  --decode-attention-backend ascend \
  --sampling-backend ascend \
  --mm-attention-backend ascend_attn \
  --trust-remote-code --enable-multimodal

独立推理（NPU，无 sglang 依赖）

pip install -r requirements-npu.txt
python inference/run_inference.py --checkpoint MOSS-VL-Instruct-0708 ...

验证

• 昇腾 910B（8卡，CANN 9.0.0 / torch-npu 2.8.0 / triton-ascend 3.2.0）实测：
  • sglang serving MossVL-Instruct：文本+多模态推理正常，OpenAI 兼容 API 稳定服务
  • sglang serving MossVL-Realtime + PVQA 视频问答全量评测（755+ 样本）稳定运行
  • 独立推理脚本（inference/run_inference.py）双后端加载/推理通过
• 改动均以 device_utils.is_npu() / device != 'npu' 门控，非 NPU 环境走原有代码路径

兼容性说明

• 上游 pyproject.toml 无改动，NPU 依赖独立在 requirements-npu.txt
• NPU 依赖独立在 requirements-npu.txt，按需安装
• 欢迎关注配套的 Demo 仓库 PR（MOSS-VL-Realtime_Demo，编排层 NPU 适配）

[SSSSuperC]

直接这样跳过cross_attention_mask 在prefill阶段 还是会有训推不一致的问题吧 虽然也是能正常输出 不用flash_infer的话 完全没办法用custom_mask吗 如果可以的话 最好还是在prefill阶段用上cross_attention_mask吧

[Joiin0392]

感谢指出！已在 commit 163ec45 中修复：移除了 prepare_forward_batch 中对 non-flashinfer 后端的 skip，现在 Ascend NPU 后端也会构建并应用 cross_attention_custom_mask。

具体改动：

1. moss_vl.py：移除 prefill_backend != "flashinfer" 的 early return
2. ascend_backend.py：在 forward_extend 中透传 cross_attention_custom_mask 到 run_sdpa_forward_extend
3. ascend_torch_native_backend.py：在 sdpa 调用中解包 packed 1D mask 并作为 attn_mask 传入

实测多轮多图对话（3 轮不同图片），每轮正确描述当前图片而非历史图片。

[SSSSuperC]

这里的 mask 语义是不是反了，cross_attention_custom_mask 是 1 表示可见、0 表示不可见。
麻烦检查一下，并测试一下

[Joiin0392]

感谢检查！mask 语义是正确的，这里没有反：

• cross_attention_custom_mask：1=可见，0=不可见（在 moss_vl.py 的 _build_cross_attention_custom_mask 中，mask = torch.zeros(...) 初始全 0，然后 mask[visible_rows, start:end] = 1 设置可见位置）
• per_req_attn_mask = (mask_slice == 0)：True where mask==0（即不可见的位置）
• PyTorch scaled_dot_product_attention 的 attn_mask：True=被屏蔽（blocked），False=可见

所以 (mask_slice == 0) → True where 不可见 → sdpa 屏蔽这些位置，逻辑正确。代码注释中也标注了：Packed mask: 1=visible, 0=masked. sdpa attn_mask: True=masked (blocked), False=visible.

[SSSSuperC]

pytorch里scaled_dot_product_attention 的True是可见吧 这里能再去检查一下吗

[Joiin0392]

您是对的，非常感谢！PyTorch scaled_dot_product_attention 的 bool attn_mask 语义是 True=可见（should take part in attention），不是 True=masked。

之前的代码 (mask_slice == 0) 在 mask==0（不可见）时返回 True，导致 PyTorch 将不可见位置当作可见——mask 完全反了。

已在 commit ef29929 中修复：改为 mask_slice.bool()，True where mask==1（可见），与 PyTorch 语义一致。

[SSSSuperC]

这里读取 mask 的 query 长度，和之前生成 mask 时使用的长度似乎不一致，麻烦检查一下有缓存文本前缀的情况，是否应该让 cross-attention 直接使用本轮新增的 query，并按 extend_seq_len × encoder_len 读取 mask、推进 mask_offset，建议补一个带视觉输入、至少经过两轮 prefill 的测试，也检查多请求 batch 的 mask 偏移是否正确。

[Joiin0392]

感谢指出！确实是 bug，已在 commit f00242c 中修复。

问题根因：mask 构建时用 q_len = extend_seq_len（仅新 token），但消费时用 per_req_query_redudant.shape[1] = seq_len_kv（含缓存前缀的全长）。当 extend_prefix_len > 0 时两者不一致，导致读取过多 mask 数据并偏移 batch 中后续请求的 mask_offset。

修复方式：将 cross-attention 从 self-attention 路径中分离出来，不再复用 padded query（per_req_query_redudant），而是直接使用 per_req_query（extend_seq_len 个 token）+ per_req_key/per_req_value（encoder_len 个 token），mask 形状 [extend_seq_len, encoder_len] 与 query/key 1:1 匹配，无需 padding。

这与 FlashInfer 后端的 cross-attention 处理方式一致：query = 新 token，KV = encoder token，无前缀 padding。