GLM-5.3-Flash 8 并发 507011 故障排查复盘

GLM-5.3-Flash 8 并发 507011 故障排查复盘

---

0. 一句话结论

FULL_DECODE_ONLY 图重放下，8 并发导致 decode batch 出现 padded / inactive 行；CausalConv1d 的 gather kernel 对这些空行写入槽位号 -1，传给 fla_npu 的 UPDATE 算子后，负数被当成合法数组下标做状态索引 → MTE 访存越界 → model stream 执行失败（retCode=0x91）→ context abort → host 侧抛出 507011。

加 MTP 能绕开：MTP 使每请求 token 数从 1 变成 6，该步骤不再属于 "pure decode"，直接落出 FULL_DECODE_ONLY 图的捕获范围，改走 eager，padded 行不再产生。

对应修复：vllm-ascend PR #19353（backport of #18878，commit f5826a4）。

---

1. 环境

1.1 版本栈

组件	实际	v0.30.0rc2 官方验证包	差异
vllm-ascend	0.30.rc2	0.30.0rc2	—
模型	GLM-5.3-Flash-w8a8	同	—
硬件	单机 A3（Atlas 800I A3）	同	—
TP / EP	TP16 + --enable-expert-parallel	同	—
CANN	9.2.0	9.2.0	✅
HDK	26.1.1	26.0.RC1	⚠️ 偏新
torch_npu	2.10.0.post7	2.10.0.post4	⚠️ 偏新

1.2 部署参数（原始）

export VLLM_USE_V2_MODEL_RUNNER=0
export PYTORCH_NPU_ALLOC_CONF=expandable_segments:True
export HCCL_OP_EXPANSION_MODE="AIV"
export HCCL_BUFFSIZE=1024

--host 0.0.0.0 --port 8027
--max-model-len 266240
--data-parallel-size 1
--tensor-parallel-size 16
--enable-expert-parallel
--seed 1024
--served-model-name glm
--safetensors-load-strategy prefetch
--max-num-seqs 32
--max-num-batched-tokens 8192
--trust-remote-code
--quantization ascend
--kv-cache-dtype int8
--limit-mm-per-prompt '{"image":1,"video":0}'
--gpu-memory-utilization 0.92
--compilation-config '{"cudagraph_mode":"FULL_DECODE_ONLY"}'
--additional-config '{"multistream_overlap_shared_expert":true}'
--no-enable-prefix-caching
--api-server-count 1

1.3 与官方 A3 单机基线的 4 处偏离

#	项	当前配置	官方基线	后续判定
1	--no-enable-prefix-caching	有（关前缀缓存）	应保持开启	偏离，但非根因
2	additional-config	只有 multistream_overlap_shared_expert	另有 recompute_scheduler_enable + enable_flashcomm1	偏离，非根因
3	MTP	去掉了	speculative-config: num_speculative_tokens=5, method=deepseek_mtp	★ 关键变量
4	版本	HDK 26.1.1 / torch_npu post7	HDK 26.0.RC1 / post4	次要嫌疑，未证实

---

2. 故障现象

• 8 并发报错；1 / 2 / 4 并发不报错（并发数指压测客户端的并发，非服务端参数）
• 加上 MTP 后可正常运行，去掉 MTP 即崩（开关式，非渐进式）
• 报错输出：

  RuntimeError: synchronize
  /torch_npu/csrc/core/npu/NPUStream.cpp:414
  NPU function error: aclrtSynchronizeStreamWithTimeout(stream()),
  error code is 507011, ERR00100 PTA call acl api failed
  

2.1 服务端关键指标（排除容量类假设）

Engine 000: Avg generation throughput ~253.6 tokens/s
Running: 8 reqs, Waiting: 0 reqs
NPU KV cache usage: 0.9% ~ 1.1%
Prefix cache hit rate: 0.0%

KV 只用 1%、Waiting 0 → 直接排除"KV/state 池打满、抢占、OOM"这一大类假设。

2.2 报错栈位置

multiproc_executor.py → worker_base.py → worker/worker.py
→ model_runner_v1.py → acl_graph.py → torch.npu.current_stream().synchronize()
Device:14, RankID:-1, pid=5493

栈停在 ACL 图重放的同步点，且只有单个 rank 报错。

---

3. 排查时间线与关键证据

3.1 第一阶段：认识 507011（不要被它骗）

概念	结论
507011	= ACL_ERROR_RT_MODEL_EXECUTE（模型执行失败）
性质	是结果，不是根因。device 侧已先出错，host 在 synchronize 时才抛出
InnerCode=0x7150050	= the model stream execute failed
retCode=0x91	device 侧同一含义
context is abort	CANN 熔断机制：context 内任一 stream 失败 → 整个 context 标记 abort → 之后所有同步操作无差别失败

推论：NPUEvent.cpp:95 那 9 层重复栈只是"层层等、层层失败"的级联回声，没有一层是真凶。

3.2 第二阶段：检索相似 issue

Issue/PR	场景	与本例关系
#19353	GLM5Next CausalConv1d padded decode → MTE invalid addr → 507011	★★★ 本例根因
#18878	同上，main 分支版本	#19353 的上游
#15903	GLM-5.3-Flash-w8a8 + A3 + TP16 + EP，16 rank 同崩，MoE TopK 507035	同模型同拓扑，错误码不同
#16828	A3 并发升高后 507011，MTE accesses invalid GM address	同错误码 + 同"低并发正常"特征
#8992	FULL_DECODE_ONLY 图重放 + 跨流操作 → rtMemcpy capture-mode conflict	同图模式
#16446	multistream_overlap_shared_expert + FlashComm1 叠加 → device 挂死	该开关有前科
#17569	GLM-5.3-Flash 在 --no-enable-prefix-caching 下失败	正好关了前缀缓存
#18804 / #18866	v0.30.0rc2 上 GLM-5.3-Flash 的 PD 挂死等	说明 rc2 上该模型稳定性在收尾

3.3 第三阶段：★ plog 抓到第一现场（决定性证据）

# 09:45:04 起（服务启动/图捕获阶段），反复 11 次
[ERROR] OP(5967) .../csrc/attention/kv_quant_sparse_flash_attention_vllm/op_host/...
        [GetAttrParaInfo] ... ============ tiling入口 ============

# 09:55:32.367
[ERROR] RUNTIME ProcLogicCqReport:Task run failed,
        device id=15, stream_id=47, task id=4, sqSwStatus=0x103f00ca

# 09:55:33.588  ← ★ 第一现场
[ERROR] RUNTIME device error core proc.cc:1548
        device id=15, stream id=80, task id=202, MTE error, retCode=0x91

# 故障任务类型
task_id=202, type=66[KERNEL_AIVEC]

# 模型执行失败
model execute task.cc:163 → model execute error retCode=0x91, task_id=42543

判读：

• MTE error = 内存搬运引擎访问了非法地址（越界）
• KERNEL_AIVEC = 故障 kernel 跑在 AI Vector Core（向量核）上 → 是纯搬运类算子，不是计算类算子
• stream_id=80/47 只是"哪条流上跑的 task 挂了"，流是管道，管道没坏，是管道里的包裹超尺寸

纠正了算子同事"不像算子问题、是流的问题"的判断：流只是现象载体，根因是流上跑的 gather kernel 访存越界。

3.4 第四阶段：PR #19353 确认对症（8 项全匹配）

维度	现场	PR #19353	匹配
模型	GLM-5.3-Flash-w8a8	GLM5Next（models/glm5next/ + ops/triton/kda/）	✅
分支	v0.30.0rc2	releases/v0.30.0 backport	✅
错误码	507011	"MTE invalid-address error / 507011 and killing the engine"	✅
第一现场	AIVEC + MTE error	gather kernel 对 inactive row 越界写	✅
触发条件	FULL_DECODE_ONLY	"decode graph with padded requests"	✅
并发特征	1/2/4 过，8 崩	"padded decode with changing live-row counts"	✅
机制	单 rank 挂	null_block_id = PAD_SLOT_ID = -1 被当 state 索引	✅
开关性	加 MTP 就不崩	脱离 pure decode 图	✅

---

4. 根因详解

4.1 完整故障链

不开 MTP，8 并发 decode
    ↓
每步是 (8请求 × 1token) 的 pure decode → 命中 FULL_DECODE_ONLY 图重放
    ↓
跑到第 N 步只剩 3 个请求，但图按 8 行录死 → 补 5 行零 = padded/inactive 行
    ↓
8 个请求的 conv_state 散落在状态池（不连续）
    ↓  if not conv_state.is_contiguous(): → 走 gather
gather kernel（AIVEC 上的纯搬运小算子，MTE 实际执行）
    ├─ 活跃行    → 搬到 packed 第 0/1/2 行
    └─ inactive  → 槽位号记为 -1（PAD_SLOT_ID）  ★ BUG
    ↓
交给 fla_npu 的 UPDATE 算子更新卷积状态
    └─ 传 null_block_id = -1
       → UPDATE 认为"空块过滤已禁用"，且把 -1 当作真实数组下标
    ↓
MTE 去搬第 -1 行 → 落在 buffer 外 → MTE error
    ↓
model stream 执行失败（retCode=0x91）→ context abort → host 侧 507011

4.2 为什么 8 才炸、1/2/4 不炸

补丁全在 if not conv_state.is_contiguous(): 分支内。触发需两个条件同时成立：

1. conv_state 不连续（才走 gather）
2. 存在 inactive 行

8 并发让这两个条件同时凑齐；低并发下要么状态恰好连续（直接原地更新，不调 gather），要么 padded 行未触发越界分支。

4.3 为什么加 MTP 就不崩（开关式证据）

	无 MTP	有 MTP（verify 阶段）
每请求 token 数	1	1 + 5 = 6
batch 形状	(num_reqs, 1)	(num_reqs, 6)
是否 pure decode	✅ 是	❌ 不是
是否走 decode 图	✅ 走 → 触发	❌ 不走 → 走 eager，不触发

"加 MTP 就不崩"是开关式（0/1）而非渐进式，这本身就是强证据：若为资源/容量/时序竞争问题，应是并发越高越易崩；实际是 MTP on/off 直接决定崩不崩 → 只能是代码路径分歧 → 指向"是否走 padded decode 图重放" → 正是 #19353 修的东西。

4.4 排除的假设（含我自己的两次误判）

假设	证伪依据
KV/state 池打满、抢占、OOM	KV 使用率仅 0.9%~1.1%，Waiting 0
--no-enable-prefix-caching 导致混合 KV 越界	服务正常跑 10 分钟、吞吐 253.6 tok/s，缓存路径是通的
★ 误判1：kv_quant_sparse_flash_attention + --kv-cache-dtype int8 是根因	那 11 条 tiling入口 是 op_host 把调试打印写成了 [ERROR] 级别；且从服务启动 09:45 就开始刷，服务随后正常运行。真正的第一现场是 09:55 的 MTE error
★ 误判2：算子同事"是流的问题"	流只是载体。根因是流上跑的 gather kernel 访存越界，属算子侧
★ 表述修正：MTP 主因是"accept 数每步不同"	更准确的主因是每请求 token 数 1 → 6，直接不再符合 pure decode 形状假设。"accept 数变化"是附加因素

---

5. 解决方案

5.1 PR #19353 / commit f5826a4 改了什么

文件 1：vllm_ascend/ops/triton/kda/conv_state.py（+4 −2）

# 旧
tl.store(packed_indices + request, tl.where(active, request, -1))
# 新
# The final packed row is a dedicated non-negative null slot.
# AscendC UPDATE disables null-block filtering for -1.
tl.store(packed_indices + request, tl.where(active, request, REQUESTS))

文件 2：vllm_ascend/models/glm5next/ops/causal_conv1d.py（+6 −4）

if not conv_state.is_contiguous():
    requests = cache_indices.shape[0]
    state_len, dim = conv_state.shape[1:]
    # 旧：kernel_state = torch.empty((requests, state_len, dim), ...)
    #     null_block_id = PAD_SLOT_ID          # = -1 ← 致命
    # 新 ↓
    null_block_id = requests                    # 非负！指向多出的那行
    kernel_state = torch.empty((requests + 1, state_len, dim), ...)  # 多开 1 行当"垃圾桶"

文件 3/4：测试（e2e +117 行，UT +6 行）

• batch 取 [4, 8, 16] × use_graph [False, True]
• live_counts = [batch, batch-1, batch-2, batch//2+1, 0, batch] → 正是"8 人录像带里只剩 5/6/7 人在跑"，与场景完全对应

5.2 只需 f5826a4，另两个 commit 互相抵消

commit	内容	是否需要
f5826a4	核心修复	✅ 只要这个
d2e9ddc	新增 release_notes/fragments/19353.yaml	❌ 不需要
6620a4d	删除 18878.yaml + 19353.yaml	❌ 不需要

注意 f5826a4 自带 release_notes/fragments/18878.yaml（+4），cherry-pick 后会多一个无害 yaml，手工删掉即可。

5.3 ⚠️ 遗留坑：kernel_state[-1].zero_() 未落地

逐行比对 diff 确认：新增行只有 torch.empty((requests + 1, ...))，没有 kernel_state[-1].zero_()。

gemini 提的 high 级意见（把预留的"垃圾桶行"清零）在当前 commit 里没有实现，尽管评论标记为 Resolved。

风险：torch.empty 分配的是未初始化显存，inactive 行会读到垃圾值/NaN → 故障可能从"崩溃"变成"静默数值错误"（后者更危险，因为不报错）。

建议手工补一行（零风险、零性能影响）：

kernel_state = torch.empty((requests + 1, state_len, dim),
                           dtype=conv_state.dtype, device=conv_state.device)
kernel_state[-1].zero_()      # ← 补上

5.4 backport 未经硬件验证（重要）

PR 原文明确：

"No separate NPU run was performed on this branch; branch CI is the validation vehicle."

main 上的 #18878 在 Ascend 910 上验证过（64+22 单测/e2e 通过、573/573 请求健康），但 releases/v0.30.0 这一支只在 CI 跑过。环境是 A3 + TP16/EP + w8a8 + int8 KV，比验证环境更严苛 → 打完补丁必须自己压一轮。

→ 验证通过后，建议把结果回帖到 #19353，这是该 backport 唯一缺的真机验证，对加速合入价值很大。

---

6. 知识沉淀

6.1 术语速查

术语	含义
507011	ACL_ERROR_RT_MODEL_EXECUTE，模型执行失败。是结果不是根因
MTE	Memory Transfer Engine，NPU 上专管搬数据的硬件单元（DMA）。负责校验地址合法性，越界即报 MTE error
AIVEC	AI Vector Core，向量核。纯搬运算子（如 gather）跑在这上面，计算算子跑在 AI Core
context abort	CANN 熔断：context 内任一 stream 失败 → 整个 context 标记 abort → 之后所有同步无差别失败
retCode=0x91	the model stream execute failed，device 侧对 507011 的表述
FULL_DECODE_ONLY	只把 pure decode（每请求 1 token）步骤录成图重放
padded / inactive 行	图为固定 8 行，实际只剩 3 个请求 → 补 5 行零。是矩阵里的行，不是进程/容器（跟 k8s pod 无关）
gather kernel	按索引把分散数据收集成连续块的纯搬运算子
conv_state	CausalConv1d 的卷积"记忆"，每个请求一份，散落在状态池
fla_npu UPDATE	flash-linear-attention 算子库里做状态更新的算子。其 null_block_id 若为负数，被当作"功能禁用"哨兵 + 真实下标
backport	把 main 上的修复移植回发布分支（本例 main #18878 → releases/v0.30.0 #19353）
plog	~/ascend/log/debug/plog/*.log，按 PID 分文件的进程级日志。⊂ CANN 日志

6.2 三种 "gather" 辨析（易混）

	MoE gating	FlashAttention 的 KV gather	CausalConv1d 的 conv_state gather
本质	计算	数据搬移	数据搬移
干什么	小线性层+softmax/topk，决定 token 去哪个专家	把分页 KV 收成连续 buffer	把卷积状态收成连续 buffer
是否必须	MoE 架构天生需要	总是做（PagedAttention 分页是常态）	条件触发：if not conv_state.is_contiguous()
目的	决策	适配分页内存布局	图重放时让地址固定

GLM-5.3-Flash 是混合架构，两条 gather 路径并存：

• 稀疏/全注意力层（MLA / FA 路径）→ KV cache gather ← 对应 plog 里的 kv_quant_sparse_flash_attention_vllm
• KDA 线性注意力层（34 层）→ CausalConv1d conv_state gather ← 本次 bug 在这

6.3 MTP 原理（澄清常见误解）

误区："MTP 让加载权重变快" → ❌ 恰恰相反，MTP 要额外加载一份 draft 权重。

真相：decode 是 memory-bound（访存瓶颈）——一次 forward 绝大部分时间花在"把几十 GB 权重从 HBM 搬到计算核心"，真正在算的占比很小。

公交车类比：车上装满货物（权重），普通 decode 每次只运 1 个乘客；MTP 一次运 6 个。油钱（搬运开销）几乎一样，但运了 6 个人。

→ MTP 不是让每次搬运变快，而是让"搬一次权重"产出更多 token（步数变少）。

两阶段流程：

1. DRAFT ：轻量 draft head 快速猜 5 个 token（成本极低）
2. VERIFY：把 1 真实 + 5 草稿 = 6 个 token 一起塞进主模型，一次 forward 验证
3. 拒绝采样：接受前 k 个，第 k+1 个用主模型自己的结果

"最后还是只吐一个 token 吗？" → 不是。一次 verify 可能产出 1~6 个 token。对外（流式）表现为 token 一串串往外冒。

6.4 方法论沉淀

① 507011 类错误的黄金法则

507011 / context is abort / 重复 synchronize 栈 = 回声，不是第一现场。 永远往前翻，找时间最早的那条 ERROR，而不是位置最靠后的那条。

② 区分"回声"与"第一现场"

• 回声：只有 model stream execute failed / stream synchronize / context is abort
• 第一现场：含 MTE / invalid GM / vector core / fftsplus / rtMemcpy / capture mode / HcommProxy / halMemMap

③ 用"开关式 vs 渐进式"快速定性

• 开关式（某参数 on/off 直接决定崩不崩）→ 代码路径分歧
• 渐进式（压力越大越易崩）→ 资源/容量/时序竞争

④ 对照实验一次只改一个变量

那次实验同时改了 4 处（图模式、batched tokens、前缀缓存、端口），结果无法归因。

⑤ 先排除再深挖，用指标证伪

KV 1% + Waiting 0 一条就证伪了整个"容量不足"大类假设。

⑥ 打补丁后必须验证"是否真的生效"

查三件事：改的是不是被加载的那份 / 两个文件是否都改了 / 分支是否真的进了。

⑦ 修复可能把"崩溃"变成"静默错误"

未初始化内存（torch.empty）比崩溃更危险。补 zero_()。

---

7. 排障命令速查

7.1 plog 采集流程

# ① 复现前设级别（否则抓不到流/任务级信息）
export ASCEND_GLOBAL_LOG_LEVEL=0        # 0=DEBUG
export ASCEND_SLOG_PRINT_TO_STDOUT=0
export HCCL_CONNECT_TIMEOUT=1200
export HCCL_EXEC_TIMEOUT=1200

# ② 停服务 → 备份 → 清空（改名而非删除）
pkill -9 -f "vllm serve"; pkill -9 -f "VLLM::EngineCore"; sleep 5
BAK=/root/cann_log_bak_$(date +%m%d_%H%M); mkdir -p $BAK
cp -r ~/ascend/log $BAK/
mv ~/ascend/log ~/ascend/log_old_$(date +%m%d_%H%M)
mkdir -p ~/ascend/log/debug/plog

# ③ 复现（只跑到崩溃就停，别继续压以免冲掉第一现场）

# ④ 崩溃后立刻抓（先别重启）
cd ~/ascend/log/debug/plog
ls -lt | head                          # 确认是本次新生成
grep -H "\[ERROR\]" *.log | awk '{ if (match($0, /[0-9]{4}-[0-9]{2}-[0-9]{2}-[0-9]{2}:[0-9]{2}:[0-9]{2}/)) { k=substr($0,RSTART,RLENGTH); print k"\t"$0 } else { print "9999\t"$0 } }' | sort | cut -f2- | head -20
# 上面这条：grep捞ERROR → awk把行内时间戳搬到行首 → sort按时间排 → cut删掉临时列 → head取最早20条

cp -r ~/ascend/log /root/cann_log_crash_$(date +%m%d_%H%M)/
dmesg -T | grep -iE "ascend|devmm|halMemMap|Device fault|Reserve addr"

7.2 常用过滤

# 按文件统计 ERROR 数（找第一现场所在 rank）
grep -c "\[ERROR\]" *.log | sort -t: -k2 -rn | head

# 找第一现场关键词
grep -lE "MTE|invalid GM|fftsplus|vector core|rtMemcpy|capture mode|HcommProxy|halMemMap" *.log

# 错误类型去重分布（快速定性）
grep -h "\[ERROR\]" *.log | sed 's/.*\[ERROR\]//' | sort | uniq -c | sort -rn | head -20

# 注意：不要用 cat *.log | grep —— 会丢掉文件名，16 个 rank 混在一起无法定位

8.3 判"慢"还是"卡"

grep -E "Avg generation throughput|Engine 000" glm5.3-flash.log | tail -5
npu-smi info -t usagemem
npu-smi info -t health -i <device_id>

• 有稳定 tok/s 且利用率 >50% → 是真的慢
• 无输出且利用率 0% → 是卡住（另一个问题，需单独查）

---

附：参考链接

• vllm-ascend PR #19353（releases/v0.30.0 backport，commit f5826a4）
• vllm-ascend PR #18878（main 分支原始修复）
• 相关 issue：#15903 / #16828 / #8992 / #16446 / #17569 / #17565 / #18804 / #18866
• vllm-ascend 官方文档：GLM-5.3-Flash A3 单机部署基线
• Ascend 官方 FAQ：507011 第一条判断即"固件驱动与 CANN 版本不配套"