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 GateSK 故障隔离是一个聚焦的后续流程follow-up前提是某个具名 SuperKernel scope 已被选定、且候选在运行时或正确性校验中失败。它的目标只有一个从运行时证据中识别出失败的子算子并产出一个排除了该子算子的依赖安全dependency-safe的下一 scope。文档明确要求把这条工作流与常规的 scope 设计、融合深度分析、性能调优流程分离避免混线诊断。在动手改任何标记marker之前必须先通过进入门控确认以下四项全部成立已知的失败候选包含确切命令、rank/device、进程 IDPID、时间戳plog与全新的sk_meta/pid产物确实属于这一次尝试attempt失败类型属于执行失败、正确性失败、超时、挂起或设备故障之一。纯编译失败或浅融合/无融合结果不属于本技能应回到父级适配/调试流程处理当前生效的 wrapper 与源码修订版本已确认到足以解读诊断字段和 scope API 的程度。若缺少任一新鲜日志或产物正确的动作是停下并报告证据缺口evidence gap而不是凭旧运行结果或仅凭生成的 SK hash 去推断失败子算子。这条红线贯穿整个工作流。二、第一步冻结并盘点证据在重跑任何东西之前先复制或记录所有不可变路径建立证据索引。技能文档给出的必填项如下表项目必备值运行身份命令、rank、device、PID、时间戳候选Candidatescope 名、源码标记位置、options、模型/配置修订版本运行时日志plog路径与故障行偏移元数据sk_meta/pid路径、sk_fused_nodes.log、sk_node_detail.log、sk_scope_split.log调度映射需要opId或scopeId时的sk_task_queue.json诊断模式执行 trace 或跨核检查是否改变了本次运行这些元数据文件在 AOT 源码中都有明确的落点可以对照确认文件名与生成时机sk_fused_nodes.log由 sk_optimizer.cpp 中的PrintSKNodes写入按 scope 打印每个 SK 函数名、scope id 及过滤后节点列表sk_node_detail.log节点明细日志见 sk_graph.cppsk_scope_split.logscope 拆分过程的日志见 sk_scope_split.cppsk_task_queue.jsonAIC/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.ascextern 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 入口信息不足以把失败归因到某个子算子。四、第三步把故障映射到子算子证据使用有严格的优先级顺序按序取用来自debug_op_exec_trace或当前版本等效诊断选项的显式子算子 start/finish 记录PrintDfxInfo给出的 PC 到符号的匹配结果包括子算子函数符号、kernel 类型与 stream当符号重复或 PC 不唯一时用opId/scopeId与sk_task_queue.json做关联用sk_fused_nodes.log中的节点 ID 与有序子符号、节点明细、以及sk_scope_split.log中第一个可操作触发点做交叉核对。归因前先读当前生效的 SK 源码或 TorchAir SuperKernel 参考确认每个被打印字段的确切含义——被调试版本的源码级错误打印才是字段语义的权威。文档同时列出了三条典型误归因陷阱不要因为它是最后一条日志就选它不要选文本上最近的名字不要因为某个 scope 边界恰好出现在故障附近就指向它。五、第四步置信度分级对归因结果必须显式记录high/medium/low三档置信度high有子算子执行记录或唯一的 PC/符号匹配且与任务队列、融合节点元数据互相一致mediumopId/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或干脆放弃该候选——绝不为了挽回融合深度而绕过排序约束。七、第六步把排除当作新候选重新验证排除方案落地后它就是一个新的候选需要按固定顺序重跑各道门禁preflight 与 option 兼容性执行与正确性覆盖全部 rank 以及此前失败的路径全新的 SK 元数据证明预期的子算子已融合、且被排除的子算子没有重新进入 SK仅当需要确认子算子顺序时才启用诊断 profiler前四项全部通过后才做干净的重复性能运行。反过来说以下信号不能证明排除是正确的编译成功、某条日志行消失、融合子算子数量变多。八、输出契约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不得删除既有的失败现场。九、失败路由表与父级交接技能文档用一张路由表覆盖了隔离过程中常见的分支情形可以直接作为运行手册症状处置无新鲜plog或sk_meta停止报告证据缺失只知道sk_entry_*继续做子算子关联但不得宣称完成归因符号重复或 PC 歧义用任务队列关联opId/scopeId或缩小 scope故障只在调试模式下消失把调试运行视为定位证据然后重新干净测试排除违反 event/barrier/依赖边在边界处拆分或放弃候选子算子已识别但无法安全分离优先保正确性报告候选被阻塞排除后候选通过重建全新元数据与干净门禁不复用过期证据最后是与父级技能的职责分界当本技能由 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创作声明：本文部分内容由AI辅助生成（AIGC），仅供参考