CANN graph-autofusion 的 SK 算子代码生成sk-operator-codegen源码形态识别、SK bind 适配与 compare 工程生成指南【免费下载链接】graph-autofusionGraph-autofusion 是一个面向昇腾Ascend芯片的轻量级、解耦式组件集合旨在通过自动融合技术加速模型执行。 目前已开源 SuperKernel 组件和 Autofuse 组件未来将持续开放更多自动融合相关模块。项目地址: https://gitcode.com/cann/graph-autofusion导读sk-operator-codegen是 CANN graph-autofusion 仓库中 SuperKernel 算子接入链路的核心代码生成组件它负责识别算子源码属于普通__global__kernel、当前 SK bind 还是无法自动处理的形态将普通 AscendC 算子自动适配为 SuperKernel 入口Args struct 模板化__sk__函数 SK_BIND并完成聚合目录与 standalone compare 工程生成。读完本文你将掌握该工具的完整命令集、--io-contract语义契约的编写规则、多 NPU arch 原生 wheel 的构建策略以及它在端到端流水线sk-operator-pipeline中的定位与单独调试方法。工具定位为什么需要单独使用sk-operator-codegen在 graph-autofusion 的 SuperKernel 接入链路中端到端场景优先使用sk-operator-pipeline run-sk-pipeline走完整闭环而sk-operator-codegen专注于生成阶段本身适合在以下场景单独定位问题见 sk-operator-codegen/README.md只想确认一个算子能否被识别生成阶段失败需要单独重跑adapt-sk-from-global想检查聚合后的目录是否满足后续 pybind/wheel 构建要求需要用apply-remediation对静态检查结果做自动修复尝试。从源码结构看该 skill 的交付边界十分清晰见 SKILL.md输出交给sk-operator-validate执行 contract/spec/compat 规则包并输出统一 findings、sk-operator-build-package消费operator-sk-adapted.json与operator-sk-adapted/生成 pybind binding、wheel 与 standalone compare 工程以及sk-operator-pipeline编排完整闭环。因此本文介绍的命令产物大多会成为下游 skill 的输入。支持的源码形态与识别结果工具首先对输入资产做形态分类CLI 入口统一为python3 skills_root/sk-operator-codegen/scripts/operator_codegen.py subcommand ...其中skills_root指向本仓库 skills 目录即.claude/skills。三种输入形态及其处理结果如下表输入形态典型特征处理结果当前 SK bind已包含当前框架需要的 bind 结构复用并生成 manifest按字节复制标记为already_current普通__global__kernel只有 device kernel 或 host wrapper 不完整生成 SK bind 封装无法识别缺少入口、规格或语义信息生成诊断结果交由人工处理detect-sk-form子命令会把每个源文件分类为none、current-sk-bind、partial或unknown整体形态还可能是mixed结果写入operator-sk-form-analysis.json对应实现见 operator_codegen.py 中的cmd_detect_sk_form。识别单个算子形态的命令python3 skills_root/sk-operator-codegen/scripts/operator_codegen.py detect-sk-form \ $OPERATOR_ASSET \ --output-dir build/examples/sk-codegen/detect$OPERATOR_ASSET可以是算子目录或单个.asc/.cpp源文件输出目录下会生成operator-sk-form-analysis.json其中包含overall_form与逐文件的form、has_global、has_sk_keyword、has_sk_bind等字段。若目录下没有任何 kernel 源文件命令会直接报错退出。核心命令一adapt-sk-from-global把普通 kernel 适配为 SK bind基本用法python3 skills_root/sk-operator-codegen/scripts/operator_codegen.py adapt-sk-from-global \ $OPERATOR_ASSET \ --output-dir build/examples/sk-codegen/adapted/my_op \ --io-contract operator-io-contract.jsonadapt-sk-from-global统一处理已分类的输入形态普通或可修复的非 SK 源码会先走 codegen 拥有的预适配自动修复再生成当前 SK bind模板化 global 会保留模板参数并追加splitidx只有字段类型依赖模板参数时Args struct 才模板化详见 sk-adaptation-cookbook.md。生成内容与 aclgraph-canonical 布局对干净的普通__global__kernel工具会在原始函数之后按顺序生成三部分且原始__global__函数保持不变Args struct命名为NameCamelArgs每个 kernel 参数对应一个字段保持原始顺序小于 4 字节的 C 类型int8_t、uint8_t、int16_t、uint16_t、bool使用alignas(4)。模板化__sk__函数templateuint32_t splitidx __sk__ kernel_type void name_sk(const NameCamelArgs *args [, sk::SkSystemArgs *sysArgs]) { c_type param args-param; // one line per parameter // ... original body verbatim ... }原始 body 默认不改动只有 body 引用了AscendC::GetBlockNum()时才注入sysArgs参数并把调用改写为sysArgs-skNumBlocks。SK_BIND语句SK_BIND(orig, mask, name_sk0, name_sk1, name_sk2, name_sk3)mask默认是 4DCCI允许值 0..7其中 0 表示没有能力 bit1/2/4 是 bit flag--num-splits控制绑定多少个name_skN符号范围 1..4。对应 CLI 参数为--mask默认4与--num-splits默认4源码中通过adapt_source_text(..., maskint(args.mask), num_splitsint(args.num_splits), ...)传入见 operator_codegen.py。生成包遵循 aclgraph-canonical 布局operator-sk-adapted/ csrc/op.asc csrc/pybind11.asc op_extension/__init__.py op_extension/_arch_selector.py op_extension/_torch_library.py setup.py此外还会生成operator-sk-adapted.jsonmanifest记录statuscompleted或needs-human、pybind_layout、package_name、per_file、escalations等元数据。Kernel 类型映射原 kernel 的 qualifier 到 SK qualifier 的映射规则见 sk-adaptation-cookbook.md原始 qualifierSK qualifier__vector____vector____cube____cube____mix__(c, v)general__mix__(c, v)__mix__(1, 0)__cube__特殊情况__mix__(0, 1)__vector__特殊情况bare__aicore____aicore__--with-sys-args注入规则--with-sys-argsauto是默认值只有原始 body 包含AscendC::GetBlockNum()时才注入sk::SkSystemArgs *sysArgs参数--with-sys-argsalways/never可以强制选择。注入后使用当前 API 名sysArgs-skNumBlocks/sysArgs-SkGetNumBlocks()——历史名称skBlockNum/SkGetBlockNum在当前 CANN 头文件下会编译失败sk-operator-validate --rule-pack spec会将其标记为sk.sys-args-api-current并支持自动重命名修复。--io-contract算子 IO 语义契约固化脚本不会根据变量名猜测输入输出这是该工具最重要的设计原则。当 kernel 有多个 tensor-like 参数时必须由用户、adapter skill 或上游资产契约明确说明每个 tensor-like 参数属于inputs、outputs还是workspaces以及 pybind 单返回值应该返回哪个 tensor。最小格式{ schema_version: 1, entries: { add_custom: { inputs: [x, y], outputs: [z], pybind_return_tensor: z } } }契约解析在源码中由_read_pybind_io_contract完成见 operator_codegen.py并包含严格校验schema_version必须为 1entries必须是对象字段名还兼容input_tensors/output_tensors/workspace_tensors/params/return_tensor等别名。典型失败诊断码工具会把无法自动处理的情况输出为needs-human并在operator-sk-adapted.json的escalations中给出带codegen.*前缀的诊断码。以下是文档与源码共同确认的常见场景codegen.pybind-return-tensor-unresolvedentry 有多个 tensor-like 参数但没有匹配的--io-contractcodegen.io-contract-tensor-incomplete契约匹配但遗漏了某个 tensor-like 参数未归入 inputs/outputs/workspacescodegen.runtime-parameter-contract-required存在 host-struct 运行时参数但未在parameters中声明codegen.io-contract-parameter-invalid/codegen.io-contract-parameter-kind-mismatch契约声明了源码中不存在的参数或声明 kind 与源码类型检测结果不兼容codegen.tensor-list-runtime-wrapper-required与codegen.tensor-list-prepared-runtime-state-required声明了 TensorList 运行时参数但缺少 graph-safe 的 prepared descriptor 策略codegen.tensor-list-descriptor-order-mismatchdescriptor_order与契约中 TensorList 参数声明顺序不一致codegen.unknown-sk-formpartial/unknown形态不猜测直接输出人工处理项codegen.io-contract-entry-unusedIO 契约没有匹配到任何已适配的 entrycodegen.pybind-entry-specialization-ambiguous同一 public entry 映射到多个 bind target 且无法消歧。struct-valued 运行时参数与地址参数struct-valued 运行时参数需要在parameters中声明例如tiling: {kind: host_struct}而GM_ADDR workspace或GM_ADDR tiling这类地址参数仍应放入workspaces不要声明成 host struct。契约校验支持的 kind 包括tensor、tensor_list、scalar、host_struct且支持nullable布尔字段launch对象可声明正整数block_dimcompile对象可声明defines自动补-D前缀与options。契约还可以携带bind_target、public_entry_name、source_entry_name、runtime_wrapper等字段用于模板化 kernel 特化、名称消歧与 TensorList 运行时包装。图捕获前准备状态prepared runtime state规则有些算子需要在图捕获前准备运行状态例如 TensorList descriptor、workspace tail 元数据或持久缓存。该 skill 的生成规则是见 README.md 与 SKILL.md 的运行时状态规则这些状态必须来自显式 contract不从源码变量名或参数顺序推断runtime wrapper 只能消费已经准备好的状态不允许在 forward/capture 路径中生成隐藏 helper kernel 来临时准备状态descriptor 顺序必须和 contract 中声明的 TensorList 参数顺序一致不一致时返回needs-human或生成错误。这条规则适用于所有需要 prepared runtime state 的算子不是某个样例的特殊逻辑。在契约侧对应的声明方式是runtime_wrapper对象source资产内相对路径、entryC 标识符、tensor_list_descriptor_strategy当前仅支持prepared_workspace_tail、prepare_entry与正整数descriptor_bytes、descriptor_order。工具会校验 wrapper 源文件存在、wrapper entry 存在于源码中且descriptor_order与 TensorList 参数声明顺序完全一致缺失或多余都会报codegen.tensor-list-descriptor-order-mismatch。核心命令二aggregate-sk-adapted聚合多个算子python3 skills_root/sk-operator-codegen/scripts/operator_codegen.py aggregate-sk-adapted \ --adapted-output-dir build/examples/sk-codegen/adapted/op_a \ --adapted-output-dir build/examples/sk-codegen/adapted/op_b \ --output-dir build/examples/sk-codegen/aggregate \ --aggregate-wheel-name op_extension \ --package-version 0.1.0该命令把多个 aclgraph-canonical 适配输出合并成一个聚合包树pybind11.asc和_torch_library.py会注册所有算子 entrycsrc/op.asc全部保留聚合内 entry 名必须唯一wheel native module 仍按 entry 和 NPU arch 拆分。聚合setup.py保持 Python import 包名为op_extension同时用用户指定的 distribution name 与 version 生成 wheel 文件名。注意--output-dir必须是空目录或不存在命令输出 JSON 摘要status、package_name、package_version、entries数。生成的 pybind 层为每个算子暴露面向用户的 bind target entry见 sk-adaptation-cookbook.md函数作用run_op通过torch.library注册的 SK-facing bind targetdifferential validation 在 baseline 和 SK context 下复用同一个入口核心命令三generate-standalone-compare生成独立对比工程python3 skills_root/sk-operator-codegen/scripts/operator_codegen.py generate-standalone-compare \ build/examples/sk-codegen/aggregate \ --output-dir build/examples/sk-codegen/standalone \ --target-chip ascend-910b \ --npu-arch dav-2201该命令生成operator-sk-standalone-verify/包含runtime_compare.asc、CMakeLists.txt、复制后的 adapted csrc 和operator-sk-standalone-verify.json。生成源码会输出launch_op_baseline和launch_op_skwrapper两者调用同一个bind_target由 runtime context 区分 baseline 与 SK。如果 fixture 声明了 device buffers/scalars会分配独立 baseline/SK buffer、回拷可比输出并输出 byte/hash 对比结果没有显式 device plan 时真实设备运行返回skipped-insufficient-runtime-spec不能伪造通过。关于 NPU arch 的硬性要求standalone compare 必须有明确的 NPU arch优先显式传--npu-arch未传时只有在--target-chip能通过官方来源映射到唯一 arch 时才继续生成可编译 CMake否则输出needs-target-arch不会静默回退到某个默认 arch。这条宁缺毋滥、绝不猜测的原则同样贯穿 arch 自动映射与运行时选择见下文。多芯片版本原生产物SK_NPU_ARCHS与运行时选择生成的 ACLGraph wheel 包支持多芯片版本原生产物规则如下见 README.md构建时优先读取SK_NPU_ARCHS可用逗号或分号传多个值例如SK_NPU_ARCHSdav-2201,dav-3510wheel 内 native module 按op x arch拆分例如op_extension.add_custom_dav_2201.so和op_extension.mul_custom_dav_3510.so可用SK_BISHENG_JOBS或流水线--jobs控制并行 bisheng 编译数两者都未设置时只使用有官方源码依据的当前环境检测。目前自动映射只覆盖Ascend950*/ torch_npu SoC enum260到dav-3510其他芯片不会静默 fallback需显式设置SK_NPU_ARCHS运行时优先读取SK_ACLGRAPH_NPU_ARCH未设置时才尝试有来源依据的 SoC 自动映射。即使 wheel 里只有一个.so也不会在无法确认目标架构时静默选择。自动修复与模板能力apply-remediation基于检查结果应用自动修复python3 skills_root/sk-operator-codegen/scripts/operator_codegen.py apply-remediation \ build/examples/sk-codegen/adapted/my_op \ build/examples/sk-codegen/spec/operator-validation-findings.jsonapply-remediation读取一个顶层包含findings列表的 JSON对asset_dir下的文件应用可自动修复项。支持四种修复 kindAUTO_REMEDIATION_KINDS定义于 sk_codegen_lib.pyrename-symbol、remove-line-containing、add-include、replace-pattern。不可自动修复项会作为人工处理项escalated输出。结果写入operator-sk-remediation.json包含applied/escalated/failed/skipped统计存在 failed 时命令以退出码 1 返回。注意该命令直接修改 asset_dir 下的文件因此应在临时副本或生成目录上运行。模板生成python3 skills_root/sk-operator-codegen/scripts/operator_codegen.py list-templates python3 skills_root/sk-operator-codegen/scripts/operator_codegen.py generate-from-template \ TEMPLATE_ID \ --param namevalue \ --output-dir build/examples/sk-codegen/template-outgenerate-from-template渲染templates/id.yaml用于闭环流水线的干净输入起点list-templates列出可用模板。仓库内置的 add_custom.yaml 展示了模板契约声明id、description、parameters如dtypekind 为choice默认float16允许[float16, float32, int32]以及files每个文件含path与 string.Template 风格的template内容。它生成一个最小的非 SK__global__ __vector__add 算子csrc/id.asc与README.md可直接作为adapt-sk-from-global的干净输入演示五 skill 闭环。生成后还会输出operator-sk-template-manifest.json。输出约定与流水线中的落盘位置生成阶段通常产出见 README.md适配后的源码目录描述算子、输入、输出、构建配置的 manifest聚合目录_aggregate供 pybind、wheel 和 standalone 阶段继续使用诊断 JSON用于说明为什么某个算子不能自动生成。在总流水线中这些文件会落到以本工具各阶段的典型输出目录为例01-detect-form/op/{inputs,outputs} 02-adapt-sk-from-global/op/{inputs,outputs} 02-adapt-sk-from-global/_aggregate/{inputs,outputs}单算子适配与聚合阶段的产物结构如下见 sk-adaptation-cookbook.md单算子适配为每个 asset 写一个 aclgraph-canonical treeoperator-sk-adapted/csrc/op.asc、csrc/pybind11.asc、op_extension/__init__.py、op_extension/_torch_library.py、setup.pyaggregate-sk-adapted消费多个这样的输出渲染包含所有csrc/op.asc、一个pybind11.asc、一个_torch_library.py和一个setup.py的聚合 tree。扩展点从 SKILL.md 的扩展点一节可以确认两个官方扩展入口新基础算子模板新增templates/id.yaml新自动修复规则扩展scripts/sk_codegen_lib.py中的AUTO_REMEDIATION_KINDS并在apply_remediation中增加处理分支。验证与使用建议python3 skills_root/sk-operator-codegen/scripts/operator_codegen.py --help运行--help可查看全部子命令及其参数。需要说明的是build_parser中对sk_codegen_lib采用懒加载见 operator_codegen.py 中注释因此即使可选依赖如 pyyaml不可用--help也能正常工作。结合上文实际使用时的几条建议先识别再适配先跑detect-sk-form确认形态current-sk-bind源码会被按字节复用无需重复适配。多 tensor 参数必须带契约不要依赖变量名猜测契约缺失或参数遗漏都会返回needs-human及对应codegen.*诊断码这正是该工具刻意设计的不猜测语义。arch 必须显式构建用SK_NPU_ARCHS、运行时用SK_ACLGRAPH_NPU_ARCHstandalone compare 用--npu-arch无法从官方来源确认时工具会拒绝生成而非静默 fallback。保持 graph capture 语义任何需要 capture 前准备的 runtime stateTensorList descriptor、workspace tail 元数据等都必须由显式契约表达并由 runtime wrapper 消费严禁在 forward 路径插入隐藏 helper kernel。总结sk-operator-codegen以识别形态 → 契约驱动适配 → 聚合 → 独立对比工程为主线把普通 AscendC__global__kernel 系统化地接入 SuperKernel它以 aclgraph-canonical 布局产出可被下游sk-operator-validate、sk-operator-build-package与sk-operator-pipeline直接消费的产物同时通过--io-contract、prepared runtime state 规则与严格的 arch 解析策略把需要人工判断的语义边界明确暴露为needs-human诊断而非靠猜测掩盖问题。对于希望单独调试生成阶段、或手工检查算子接入质量的开发者本文的命令集与契约格式可以直接照搬使用。延伸阅读sk-operator-codegen 主 README本工具命令与输出约定的权威说明sk-operator-codegen SKILL.md命令速查、扩展点与运行时状态规则SK 适配规则手册Args struct /__sk__/SK_BIND生成细节与 kernel 类型映射operator_codegen.pyCLI 全部子命令与 IO 契约解析实现sk_codegen_lib.py自动修复规则与代码生成库实现add_custom.yaml内置最小非 SK 算子模板示例sk-operator-pipeline 目录端到端闭环编排入口sk-operator-validate 目录contract/spec/compat 规则包【免费下载链接】graph-autofusionGraph-autofusion 是一个面向昇腾Ascend芯片的轻量级、解耦式组件集合旨在通过自动融合技术加速模型执行。 目前已开源 SuperKernel 组件和 Autofuse 组件未来将持续开放更多自动融合相关模块。项目地址: https://gitcode.com/cann/graph-autofusion创作声明：本文部分内容由AI辅助生成（AIGC），仅供参考