graph-autofusion SuperKernel SK 故障隔离：从 plog 现场定位失败子算子并生成安全排除候选

graph-autofusion SuperKernel SK 故障隔离：从 plog 现场定位失败子算子并生成安全排除候选

【免费下载链接】graph-autofusionGraph-autofusion 是一个面向昇腾（Ascend）芯片的轻量级、解耦式组件集合，旨在通过自动融合技术加速模型执行。 目前已开源 SuperKernel 组件和 Autofuse 组件，未来将持续开放更多自动融合相关模块。项目地址: https://gitcode.com/cann/graph-autofusion

本文基于 graph-autofusion 仓库中 SuperKernel 调试技能文档 superkernel-sk-failure-isolation，讲解当一个已被选定的 SuperKernel scope 候选在执行、正确性、超时或挂起阶段失败时，如何以全新plog与sk_meta证据为主线，把故障精确定位到具体子算子，并生成依赖安全的下一轮 scope 排除方案。读完后，你能掌握完整的 SK 故障隔离工作流：证据封存清单、plog 故障标记检索、子算子归因的四级证据链、置信度分级、排除提案要素，以及配套的失败现场输出契约与父级技能交接规则。

一、技能定位与进入门控（Entry Gate）

SK 故障隔离是一个聚焦的后续流程（follow-up），前提是"某个具名 SuperKernel scope 已被选定、且候选在运行时或正确性校验中失败"。它的目标只有一个：从运行时证据中识别出失败的子算子，并产出一个排除了该子算子的依赖安全（dependency-safe）的下一 scope。文档明确要求把这条工作流与常规的 scope 设计、融合深度分析、性能调优流程分离，避免混线诊断。

在动手改任何标记（marker）之前，必须先通过进入门控，确认以下四项全部成立：

• 已知的失败候选包含：确切命令、rank/device、进程 ID（PID）、时间戳；
• plog与全新的sk_meta/<pid>产物确实属于这一次尝试（attempt）；
• 失败类型属于执行失败、正确性失败、超时、挂起或设备故障之一。纯编译失败或"浅融合/无融合"结果不属于本技能，应回到父级适配/调试流程处理；
• 当前生效的 wrapper 与源码修订版本已确认到足以解读诊断字段和 scope API 的程度。

若缺少任一新鲜日志或产物，正确的动作是停下并报告证据缺口（evidence gap），而不是凭旧运行结果或仅凭生成的 SK hash 去推断失败子算子。这条红线贯穿整个工作流。

二、第一步：冻结并盘点证据

在重跑任何东西之前，先复制或记录所有不可变路径，建立"证据索引"。技能文档给出的必填项如下表：

这些元数据文件在 AOT 源码中都有明确的落点，可以对照确认文件名与生成时机：

• sk_fused_nodes.log：由 sk_optimizer.cpp 中的PrintSKNodes写入，按 scope 打印每个 SK 函数名、scope id 及过滤后节点列表；
• sk_node_detail.log：节点明细日志，见 sk_graph.cpp；
• sk_scope_split.log：scope 拆分过程的日志，见 sk_scope_split.cpp；
• sk_task_queue.json：AIC/AIV 任务队列的 JSON 导出，生成逻辑见 sk_dump_json.cpp。

这里有一条关键的诊断纪律：诊断性运行（开启调试选项）与干净的性能够据必须分开保存。因为调试选项本身可能改变融合结果与运行计时，把二者混在一起会得到无法复现、不可比的证据。

三、第二步：在 plog 中定位 SK 故障

对新鲜的plog，按当前版本的故障标记检索。技能文档给出的已知标记集如下（并建议补充本地源码实际打印的精确字符串）：

fault kernel_name Exception is from superkernel function PrintDfxInfo sk_entry_aic sk_entry_aiv sk_entry_mix opId scopeId COND PC

这些标记不是凭空约定的，它们直接对应本仓库 SuperKernel 异常处理器的行为。在 sk_dfx_exception_handler.cpp 中可以看到判定逻辑：

// Check if function name starts with sk_entry (e.g., sk_entry, sk_entry_aiv, sk_entry_mix11, etc.) if (!StartsWith(funcName, "sk_entry")) { SK_DLOGD("fault kernel_name '%s' does not start with 'sk_entry', skipping", funcName); return false; } ... SK_DLOGE("Exception is from superkernel function '%s', op_trace=%s, proceeding with handling", funcName, hasOpTrace_ ? "true" : "false");

也就是说，fault kernel_name ... does not start with 'sk_entry'这条日志意味着该设备故障不属于 SK kernel，应跳过 SK 归因流程；而Exception is from superkernel function出现时，处理器才继续从异常信息中提取设备侧参数、拷贝任务队列、解析 PC/COND 寄存器并调用PrintDfxInfo（见 sk_dfx_exception_handler.cpp）。异常处理器随后会把故障 PC 与每个子算子的入口地址区间做匹配（FindErrorNodeByPC/IdentifyErrorNodeByPC，见 sk_dfx_exception_handler.h），这正是"PC→子算子符号"这条证据链的源码实现。

而sk_entry_aic、sk_entry_aiv、sk_entry_mix等入口名则直接来自设备端入口 kernel 源码 sk_entry.asc：

extern "C" __global__ __attribute__((aligned(512))) __mix__(0, 1) void sk_entry_aiv(GM_ADDR skDevArgs) { ... } extern "C" __global__ __attribute__((aligned(512))) __mix__(1, 1) void sk_entry_mix11(GM_ADDR skDevArgs) { ... } extern "C" __global__ __attribute__((aligned(512))) __mix__(1, 0) void sk_entry_aic(GM_ADDR skDevArgs) { ... }

同一文件中还派生出_debug、_dump_profiling、_op_trace、_early_start等组合变体（如sk_entry_aic_op_trace），入口名中的op_trace字样会被异常处理器捕获（hasOpTrace_），触发额外的子 kernel 符号打印。

一个硬性确认点：必须确认报告的入口确实是一个 SK 入口。一次泛化的设备故障（没有 SK 入口信息）不足以把失败归因到某个子算子。

四、第三步：把故障映射到子算子

证据使用有严格的优先级顺序，按序取用：

1. 来自debug_op_exec_trace（或当前版本等效诊断选项）的显式子算子 start/finish 记录；
2. PrintDfxInfo给出的 PC 到符号的匹配结果，包括子算子函数符号、kernel 类型与 stream；
3. 当符号重复或 PC 不唯一时，用opId/scopeId与sk_task_queue.json做关联；
4. 用sk_fused_nodes.log中的节点 ID 与有序子符号、节点明细、以及sk_scope_split.log中第一个可操作触发点做交叉核对。

归因前先读当前生效的 SK 源码（或 TorchAir SuperKernel 参考），确认每个被打印字段的确切含义——被调试版本的源码级错误打印才是字段语义的权威。文档同时列出了三条典型误归因陷阱：不要因为它是最后一条日志就选它；不要选文本上最近的名字；不要因为某个 scope 边界恰好出现在故障附近就指向它。

五、第四步：置信度分级

对归因结果必须显式记录high/medium/low三档置信度：

• high：有子算子执行记录，或唯一的 PC/符号匹配，且与任务队列、融合节点元数据互相一致；
• medium：opId/scopeId与元数据一致，但子符号或 PC 是重复出现的；
• low：仅凭 SK 入口、宽泛的错误消息或不完整日志才能指认候选。

low置信度时的处置规则同样明确：不得任意排除某个子算子。正确做法是用当前 wrapper 接受的诊断选项复现，首选debug_op_exec_trace；只有当失败形态指向同步或资源行为问题时，才升级到跨核检查（cross-core）或逐算子核检查（per-op core check）。若歧义仍无法消除，就缩小 scope 范围，直到把失败子算子单独隔离出来。

六、第五步：生成下一 scope 的排除提案

排除提案必须是显式的、可执行的，包含五个要素：

• 要排除的确切子算子：符号、node/op ID、stream、scope ID；
• 原始 scope 区间，以及新的 before/after 区间；
• 当前 wrapper/源码修订版本支持的 scope API 或 marker 形式；
• 依赖、event、barrier、cache 与 stream 排序约束；
• 排除该子算子的理由及其证据置信度。

操作原则是"最小排除"：只排除被牵连的子算子和其直接不安全的依赖片段，保留其余安全子算子与原图依赖顺序。排除机制优先使用运行时明确提供的 unfusible/Nonemarker 机制；不可用时，就在该子算子前后对具名 scope 做拆分。文档特别警告：不要从其他 CANN 版本"发明" API 签名。

当被排除子算子参与了必需的 event、wait/notify、全核 barrier、集合通信（collective）、cache 变更或跨 stream 依赖时，必须保留这条边。若单点子算子排除无法表达这一安全边界，就应在依赖边界前后拆分 scope，或干脆放弃该候选——绝不为了挽回融合深度而绕过排序约束。

七、第六步：把排除当作新候选重新验证

排除方案落地后，它就是一个新的候选，需要按固定顺序重跑各道门禁：

1. preflight 与 option 兼容性；
2. 执行与正确性，覆盖全部 rank 以及此前失败的路径；
3. 全新的 SK 元数据，证明预期的子算子已融合、且被排除的子算子没有重新进入 SK；
4. 仅当需要确认子算子顺序时，才启用诊断 profiler；
5. 前四项全部通过后，才做干净的重复性能运行。

反过来说，以下信号不能证明排除是正确的：编译成功、某条日志行消失、融合子算子数量变多。

八、输出契约（Output Contract）

诊断结束时必须回传一份固定结构的记录给父级技能，且该记录必须满足父级技能的失败现场报告契约，并追加到失败实验自己的报告中——本技能的响应文本或父级最终摘要不能替代实验本地的记录：

Failed attempt: rank/device/PID, command, timestamp, artifact paths Failure phase: execution | correctness | timeout | hang | device fault Fault evidence: file:line or field, with a short exact excerpt Suspect child: operator/symbol, node or op ID, stream, scope ID Confidence: high | medium | low; why Exclusion proposal: exact child/range and marker or scope change Safety notes: dependencies, events, barriers, streams, cache state Verification: each gate result, artifact path, or not_run reason Residual risk: unresolved ambiguity or follow-up needed

配套的实验管理规则是：父级技能必须把失败候选保留为负面证据，把排除提案作为新一轮实验，而不是覆盖失败结果；成功的重试只是追加新 attempt，不得删除既有的失败现场。

九、失败路由表与父级交接

技能文档用一张路由表覆盖了隔离过程中常见的分支情形，可以直接作为运行手册：

最后是与父级技能的职责分界：当本技能由 superkernel-auto-tune 调用时，输出记录完成后即交还控制权。scope 源码编辑、实验编号管理、基线对比与中文最终报告都属于父级技能职责，本技能内不得重复这些工作。

十、小结：把"故障"变成可执行的下一轮实验

这套流程的核心思想是把 SK 运行失败从一个"黑盒报错"变成一条可审计的证据链：先用 Entry Gate 保证证据新鲜且版本可解读，再用plog中的fault kernel_name/Exception is from superkernel function/PrintDfxInfo标记确认故障确实落在 SK 入口内（对应 sk_dfx_exception_handler.cpp 的判定逻辑），然后按"执行 trace → PC/符号匹配 →opId/scopeId关联 → 元数据交叉核对"的四级证据顺序归因子算子，给出置信度，最终输出一个带安全约束的排除提案，并以新候选身份走完全部验证门禁。仓库中 sk_entry.asc 的入口 kernel 家族、sk_dump_json.cpp 的任务队列导出、sk_optimizer.cpp 的sk_fused_nodes.log输出，以及父级技能的失败现场契约与 TorchAir 参考，共同构成了这套流程可直接落地的证据基础。

【免费下载链接】graph-autofusionGraph-autofusion 是一个面向昇腾（Ascend）芯片的轻量级、解耦式组件集合，旨在通过自动融合技术加速模型执行。 目前已开源 SuperKernel 组件和 Autofuse 组件，未来将持续开放更多自动融合相关模块。项目地址: https://gitcode.com/cann/graph-autofusion