SuperKernel 算子代码生成实战用 sk-operator-codegen 将普通 AscendC kernel 适配为 SK bind【免费下载链接】graph-autofusionGraph-autofusion 是一个面向昇腾Ascend芯片的轻量级、解耦式组件集合旨在通过自动融合技术加速模型执行。 目前已开源 SuperKernel 组件和 Autofuse 组件未来将持续开放更多自动融合相关模块。项目地址: https://gitcode.com/cann/graph-autofusionSuperKernelSK是 CANN graph-autofusion 开源仓库中面向昇腾芯片的自动融合组件。要把一个普通 AscendC__global__算子接入 SuperKernel 执行框架核心工作是为它生成 SK bindingArgs struct 模板化__sk__函数 SK_BIND。本文围绕仓库内.claude/skills/sk-operator-codegen/SKILL.md定义的内置代码生成 skill完整讲解源码形态识别、SK bind 自动生成、多算子聚合、standalone 对比工程、自动修复与模板扩展的整套流程并结合operator_codegen.py、sk_codegen_lib.py等源码说明底层生成规则让你可以直接在本地复现一条「普通算子 → SK 适配 → 聚合 → 对比验证」的完整链路。一、skill 定位把「普通__global__算子」变成「SK 入口」在 SuperKernel 框架中算子需要以__sk__模板函数的形式暴露给 SK 运行时并通过SK_BIND宏注册。这一层源码通常被称为「SK bind」。本 skill 的职责就是把普通 AscendC__global__kernel 适配为 SuperKernel 入口对已经是当前 SK bind 形态的源码按字节复用不做任何改动对普通或可修复的非 SK 源码先执行 codegen 拥有的预适配自动修复再生成当前 SK bind对混合形态或信息不足的输入输出明确的人工处理结论如needs-human绝不猜测。自动生成逻辑以该 skill 目录下的三部分为稳定契约详见 SKILL.mdscripts/代码生成实现operator_codegen.pyCLI 入口、sk_codegen_lib.py生成核心库templates/可渲染的新算子模板如add_custom.yamlreferences/适配规则手册sk-adaptation-cookbook.md。所有命令统一通过一个入口脚本调用python3 skills_root/sk-operator-codegen/scripts/operator_codegen.py subcommand ...其中skills_root指包含sk-operator-codegen目录的 skill 根路径。在仓库中对应 operator_codegen.py其main()通过build_parser()构建子命令分发见源码末尾的main()与__main__入口。二、源码形态识别detect-sk-form适配前首先要回答一个问题这份算子源码到底是什么形态detect-sk-form子命令负责把源码分类并输出operator-sk-form-analysis.json。python3 skills_root/sk-operator-codegen/scripts/operator_codegen.py detect-sk-form \ $OPERATOR_ASSET \ --output-dir build/examples/sk-codegen/detect支持四类形态源码实现见 operator_codegen.py 的cmd_detect_sk_form逐文件调用sk_codegen_lib.detect_sk_form后汇总输入形态典型特征处理结果none普通__global__kernel无 SK 标记生成 Args struct、模板化__sk__和SK_BINDcurrent-sk-bind已包含当前框架需要的 bind 结构__sk__、SK_BIND、CommArgs/SkSystemArgs等按字节复制源码标记already_currentpartial部分 SK 特征存在但不完整不猜测输出codegen.unknown-sk-form等人工处理项unknown缺少入口、规格或语义信息输出诊断结果交由人工处理当多个源文件形态不一致时总览结果会被归类为mixed。detect_sk_form的标记检测逻辑见 sk_codegen_lib.py会同时检查__sk__关键字、SK_BIND宏以及 SK 参数结构体特征CommArgs、SkSystemArgs、__gm__ uint64_t *param并给出每项标记的置信度与证据文件。三、核心命令adapt-sk-from-global 自动生成 SK bind形态识别确认算子属于none普通__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.json该命令统一处理已分类的输入形态普通或可修复的非 SK 源码先走 codegen 拥有的预适配自动修复如自动补#include kernel_operator.h、清理仅__global__可用的 kernel task 宏、为TPipe补充Destroy等见 sk_codegen_lib.py 的adapt_source_text再生成当前 SK bind。模板化 global 会保留模板参数并追加splitidx只有字段类型依赖模板参数时Args struct 才模板化。3.1 生成包遵循 aclgraph-canonical 布局每个算子适配后输出一棵标准的 aclgraph-canonical 目录树operator-sk-adapted/ csrc/op.asc # 适配后的算子源码含 SK bind csrc/pybind11.asc # pybind11 绑定源码 op_extension/__init__.py op_extension/_torch_library.py setup.py # wheel 构建脚本3.2 IO contract多 tensor 参数的语义契约当 kernel 有多个 tensor-like 参数时--io-contract FILE是必选项。固化脚本不会根据变量名或参数顺序猜测输入输出必须由用户、adapter skill 或上游资产契约明确说明每个 tensor-like 参数属于inputs、outputs还是workspaces以及 pybind 单返回值应该返回哪个 tensor。最小格式如下{ schema_version: 1, entries: { add_custom: { inputs: [x, y], outputs: [z], pybind_return_tensor: z } } }契约解析逻辑在 operator_codegen.py 的_read_pybind_io_contract中实现支持以下完整字段inputs/outputs/workspacestensor-like 参数的归属声明workspaces用于GM_ADDR workspace、GM_ADDR tiling这类地址参数注意这类地址参数应放入workspaces不要声明成 host structpybind_return_tensorpybind 单返回值对应的输出 tensor必须在outputs中列出parameters.name.kind运行时参数种类支持tensor、tensor_list、scalar、host_struct四种struct-valued 运行时参数例如tiling: {kind: host_struct}必须显式声明否则契约解析会直接报错parameters.name.nullable布尔值标记参数是否可空launch.block_dim正整数声明算子 launch 的 block 维度compile.defines/compile.options编译宏与编译选项宏会自动补-D前缀runtime_wrapperTensorList 等复杂运行时参数的自定义 wrapper 声明见下文「图捕获前准备状态」。3.3 契约缺失或错误时的明确结论契约驱动意味着「宁缺毋滥」缺失或不完整时不会做 best-effort 兜底而是输出needs-human或具体 codegen finding多 tensor-like 参数但无--io-contract→ 返回needs-human在operator-sk-adapted.json中给出codegen.pybind-return-tensor-unresolved契约匹配但遗漏了某个 tensor-like 参数 →codegen.io-contract-tensor-incomplete存在 host-struct 参数但未声明 →codegen.runtime-parameter-contract-required契约声明了未知参数、参数 kind 与源码检测不一致、bind_target不匹配、TensorList 参数缺少 runtime wrapper → 分别输出codegen.io-contract-parameter-invalid、codegen.io-contract-parameter-kind-mismatch、codegen.io-contract-bind-target-mismatch、codegen.tensor-list-runtime-wrapper-required。这些 preflight 检查在cmd_adapt_sk_from_global中通过apply_io_contract逐 entry 执行见 operator_codegen.py。四、SK bind 生成形态详解对每一种支持的源码形态adapter 要么生成当前 SK binding要么返回明确的人工处理项。对于干净的none形态会在原始__global__函数之后按顺序生成三样东西规则契约见 sk-adaptation-cookbook.md代码实现见 sk_codegen_lib.py 的render_sk_adaptation。4.1 Args structkernel 参数的统一载体struct NameCamelArgs { c_type param; // 每个 kernel 参数对应一个字段保持原始顺序 // ... };命名为NameCamelArgskernel 名转 CamelCase 后加Args后缀每个 kernel 参数对应一个字段保持原始顺序C 类型小于 4 字节的字段int8_t、uint8_t、int16_t、uint16_t、bool自动加alignas(4)保证结构体对齐语义源码中_is_small_int_type判断 align_prefix拼接若原始 kernel 是模板函数只有字段类型依赖模板参数时 Args struct 才模板化只影响 body 或 kernel type 的模板参数保留在 SK 函数上不改变 runtime 参数包布局。4.2 模板化__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 默认不改动逐行解包args后按原样保留只有 body 引用了AscendC::GetBlockNum()时才注入sysArgs参数并把调用改写为sysArgs-skNumBlocks见 sk_codegen_lib.py 的正则改写逻辑模板参数追加uint32_t splitidx用于 SK 的多核拆分split标识。4.3 SK_BIND 注册语句SK_BIND(orig, mask, name_sk0, name_sk1, name_sk2, name_sk3)mask默认是4DCCI/早启能力位允许值0..70表示没有能力 bit1/2/4是 bit flag源码中校验mask in range(0, 8)--num-splits控制绑定多少个name_skN符号范围1..4源码中校验1 num_splits 4超出即抛错原始__global__函数必须保持不变。4.4 Kernel 类型映射__sk__函数会继承原始 kernel 的执行类型 qualifier映射规则如下源码见map_kernel_type_for_sk原始 qualifierSK qualifier__vector____vector____cube____cube____mix__(c, v)general__mix__(c, v)__mix__(1, 0)__cube__特殊情况__mix__(0, 1)__vector__特殊情况bare__aicore____aicore__4.5 sysArgs 注入策略--with-sys-argsauto是默认值只有原始 body 包含AscendC::GetBlockNum()时才注入sk::SkSystemArgs *sysArgs--with-sys-argsalways/never可以强制选择实现见 sk_codegen_lib.py 的sys_args_mode三分支。注入后使用当前 API 名sysArgs-skNumBlocks/sysArgs-SkGetNumBlocks()。历史写法skBlockNum/SkGetBlockNum在当前 CANN 头文件下会编译失败sk-operator-validate --rule-pack spec会将其标记为sk.sys-args-api-current并支持自动重命名修复。五、图捕获前准备状态codegen 的运行时状态规则部分算子需要在图捕获前准备运行状态例如 TensorList descriptor、workspace tail 元数据或持久缓存。本 skill 对这类场景有明确的硬性规则见 SKILL.md 的「运行时状态规则」一节tensor-like 参数角色不明确时必须要求 IO contractcapture 前必须准备的 runtime state例如 TensorList descriptor 或持久 workspace 元数据必须由显式契约表达并由 runtime wrapper 消费生成 wrapper不得在 captured forward 路径里插入隐藏 helper kernel来临时制造 descriptor 或状态descriptor 顺序必须由 contract 驱动并与声明的 TensorList 参数校验一致不一致时是生成错误不做 best-effort fallback。在 IO contract 中TensorList 参数需要声明kind tensor_list同时通过runtime_wrapper声明自定义 wrapper 源码与入口例如{ schema_version: 1, entries: { concat_custom: { inputs: [inputs_list], outputs: [out], pybind_return_tensor: out, parameters: { inputs_list: { kind: tensor_list } }, runtime_wrapper: { source: runtime_wrapper.cpp, entry: concat_wrapper, tensor_list_descriptor_strategy: prepared_workspace_tail, prepare_entry: concat_prepare, descriptor_bytes: 64, descriptor_order: [inputs_list] } } } }tensor_list_descriptor_strategy目前仅支持prepared_workspace_tail一种取值且一旦声明就必须同时提供prepare_entry、正整数descriptor_bytes与descriptor_order校验见 operator_codegen.py 的_runtime_wrapper_contract。wrapper 的source必须是 asset 内部的相对路径禁止绝对路径与..越界entry必须是合法 C 标识符。这条规则适用于所有需要 prepared runtime state 的算子不是某个样例的特殊逻辑。六、多算子聚合aggregate-sk-adapted单个算子适配后仍然各自独立如果希望把它们打包进一个wheel需要聚合步骤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 布局operator-sk-adapted/csrc/*.asc、csrc/pybind11.asc、op_extension/、setup.py但有以下变化pybind11.asc和_torch_library.py会注册所有算子的 entry聚合内 entry 名必须唯一重复会作为错误处理生成的 pybind 层为每个算子暴露面向用户的 bind target entryrun_op通过torch.library注册differential validation 在 baseline 和 SK context 下复用同一个入口聚合setup.py保持 Python import 包名为op_extension同时用用户指定的 distribution name--aggregate-wheel-name和 version--package-version生成 wheel 文件名wheel 的 native module 仍按entry x arch拆分详见下文第八节。七、生成 standalone compare 工程聚合完成后可以生成一个独立的对比验证工程用于在真实设备上对比 baseline 与 SK 的执行结果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可直接编译的 CMake 工程复制后的 adapted csrcoperator-sk-standalone-verify.json工程描述。生成的源码会输出launch_op_baseline和launch_op_sk两个 wrapper两者调用同一个bind_target由 runtime context 区分 baseline 与 SK保证对比口径一致。7.1 设备与 arch 的边界约束如果 fixture 声明了 device buffers/scalars会分配独立的baseline/SK buffer、回拷可比输出并输出byte/hash 对比结果没有显式 device plan 时真实设备运行返回skipped-insufficient-runtime-spec不能伪造通过standalone 的--npu-arch必须显式或只能由--target-chip解析到唯一的 source-backed arch否则输出needs-target-arch不会静默回退到某个默认 arch。八、多 arch 支持SK_NPU_ARCHS 与运行时 arch 选择生成的 ACLGraph wheel 支持多芯片版本原生产物arch 处理贯穿构建期与运行期构建期优先读取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 编译数自动映射SK_NPU_ARCHS与SK_BISHENG_JOBS都未设置时只使用有官方源码依据的当前环境检测。目前自动映射只覆盖Ascend950*/ torch_npu SoC enum260到dav-3510其他芯片不会静默 fallback需显式设置SK_NPU_ARCHS映射实现见 sk_codegen_lib.py 的_detect_soc_version/map_soc_version_to_arch运行期优先读取SK_ACLGRAPH_NPU_ARCH未设置时才尝试有来源依据的 SoC 自动映射。即使 wheel 里只有一个.so也不会在无法确认目标架构时静默选择。九、自动修复apply-remediation静态检查如sk-operator-validate产出的 findings 中有一部分是机器可自动修复的。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.json支持的自动修复种类对应sk_codegen_lib.py中的AUTO_REMEDIATION_KINDS常量rename-symbol符号重命名remove-line-containing删除包含指定内容的行add-include添加头文件包含replace-pattern正则模式替换。不可自动修复的项会作为人工处理项输出交由开发者处理。新增自动修复规则的方法是扩展scripts/sk_codegen_lib.py中的AUTO_REMEDIATION_KINDS并在apply_remediation中增加对应处理分支。十、模板生成从干净输入开始闭环流水线需要一个「干净输入」起点generate-from-template和list-templates提供这一能力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-out模板渲染templates/id.yaml。仓库内置的add_custom模板见 add_custom.yaml是一个最小化的非 SK elementwise add 算子生成干净的__global__ __vector__kernel可以直接喂给adapt-sk-from-global派生 SK 适配id: add_custom parameters: - name: dtype kind: choice default: float16 allowed: [float16, float32, int32] files: - path: csrc/$id.asc template: | #include kernel_operator.h // ... class KernelAdd 与 extern C __global__ __vector__ void $id(...) ...模板使用string.Template风格占位符$id、$dtype、$ctype因此 kernel 代码中的花括号无需转义。参数支持choice类型并声明allowed取值集合。新增基础算子模板的方法是新增templates/id.yaml。十一、输出约定与流水线落位生成阶段通常产出四类文件适配后的源码目录描述算子、输入、输出、构建配置的manifest聚合目录_aggregate供 pybind、wheel 和 standalone 阶段继续使用诊断 JSON用于说明为什么某个算子不能自动生成。在总流水线中这些文件会落到约定目录inputs/outputs分别承接上游输入与下游产物01-detect-sk-form/op/{inputs,outputs} 02-adapt-sk-from-global/op/{inputs,outputs} 02-adapt-sk-from-global/_aggregate/{inputs,outputs}十二、交付边界与其他 skill 的衔接本 skill 的输出交给下游三个 skill 继续消费见 SKILL.md 的「交付边界」一节sk-operator-validate执行 contract/spec/compat 规则包输出统一 findingssk-operator-build-package消费operator-sk-adapted.json和operator-sk-adapted/生成 pybind binding、wheel并构建 standalone compare 工程sk-operator-pipeline run-sk-pipeline编排完整闭环。需要说明的是端到端场景优先使用sk-operator-pipeline run-sk-pipeline本工具更适合单独定位生成阶段的问题见 README.md。何时直接运行本工具只想确认一个算子是否能被识别生成阶段失败需要单独重跑adapt-sk-from-global想检查聚合后的目录是否满足后续 pybind/wheel 构建要求需要用apply-remediation对静态检查结果做自动修复尝试。十三、历史 scaffold 入口兼容保留intake、plan、analyze-sk-conversion、adapt-sk-binding-scaffold、generate-sk-source-scaffold作为历史 scaffold 入口仍保留详细说明见references/workflow.md。新代码路径统一收敛到上文所述的各子命令。十四、验证与快速自检运行入口脚本的--help可以快速验证 skill 环境是否可用python3 skills_root/sk-operator-codegen/scripts/operator_codegen.py --help一条完整的最小验证路径是# 1. 用内置模板生成干净输入 python3 skills_root/sk-operator-codegen/scripts/operator_codegen.py generate-from-template \ add_custom --param dtypefloat16 \ --output-dir build/examples/sk-codegen/template-out # 2. 识别形态预期none python3 skills_root/sk-operator-codegen/scripts/operator_codegen.py detect-sk-form \ build/examples/sk-codegen/template-out --output-dir build/examples/sk-codegen/detect # 3. 生成 SK bind单 tensor 输出可不带 --io-contract python3 skills_root/sk-operator-codegen/scripts/operator_codegen.py adapt-sk-from-global \ build/examples/sk-codegen/template-out --output-dir build/examples/sk-codegen/adapted/add_custom总结sk-operator-codegen以「契约驱动、绝不猜测」为设计主线形态识别负责分类IO contract 负责把 tensor 语义与运行时状态显式化adapt-sk-from-global生成标准的 Args struct __sk__SK_BIND三件套aggregate-sk-adapted与generate-standalone-compare打通多算子打包与真实设备对比验证而模板与AUTO_REMEDIATION_KINDS则提供了可持续扩展的入口。对于希望把普通 AscendC 算子接入 SuperKernel 的开发者这条从detect-sk-form到generate-standalone-compare的链路就是最直接的实操路径对于框架开发者sk-adaptation-cookbook.md 中关于 Args struct 对齐、kernel 类型映射、sysArgs 注入的规则则是理解生成产物正确性的权威依据。【免费下载链接】graph-autofusionGraph-autofusion 是一个面向昇腾Ascend芯片的轻量级、解耦式组件集合旨在通过自动融合技术加速模型执行。 目前已开源 SuperKernel 组件和 Autofuse 组件未来将持续开放更多自动融合相关模块。项目地址: https://gitcode.com/cann/graph-autofusion创作声明：本文部分内容由AI辅助生成（AIGC），仅供参考
SuperKernel 算子代码生成实战：用 sk-operator-codegen 将普通 AscendC kernel 适配为 SK bind
SuperKernel 算子代码生成实战用 sk-operator-codegen 将普通 AscendC kernel 适配为 SK bind【免费下载链接】graph-autofusionGraph-autofusion 是一个面向昇腾Ascend芯片的轻量级、解耦式组件集合旨在通过自动融合技术加速模型执行。 目前已开源 SuperKernel 组件和 Autofuse 组件未来将持续开放更多自动融合相关模块。项目地址: https://gitcode.com/cann/graph-autofusionSuperKernelSK是 CANN graph-autofusion 开源仓库中面向昇腾芯片的自动融合组件。要把一个普通 AscendC__global__算子接入 SuperKernel 执行框架核心工作是为它生成 SK bindingArgs struct 模板化__sk__函数 SK_BIND。本文围绕仓库内.claude/skills/sk-operator-codegen/SKILL.md定义的内置代码生成 skill完整讲解源码形态识别、SK bind 自动生成、多算子聚合、standalone 对比工程、自动修复与模板扩展的整套流程并结合operator_codegen.py、sk_codegen_lib.py等源码说明底层生成规则让你可以直接在本地复现一条「普通算子 → SK 适配 → 聚合 → 对比验证」的完整链路。一、skill 定位把「普通__global__算子」变成「SK 入口」在 SuperKernel 框架中算子需要以__sk__模板函数的形式暴露给 SK 运行时并通过SK_BIND宏注册。这一层源码通常被称为「SK bind」。本 skill 的职责就是把普通 AscendC__global__kernel 适配为 SuperKernel 入口对已经是当前 SK bind 形态的源码按字节复用不做任何改动对普通或可修复的非 SK 源码先执行 codegen 拥有的预适配自动修复再生成当前 SK bind对混合形态或信息不足的输入输出明确的人工处理结论如needs-human绝不猜测。自动生成逻辑以该 skill 目录下的三部分为稳定契约详见 SKILL.mdscripts/代码生成实现operator_codegen.pyCLI 入口、sk_codegen_lib.py生成核心库templates/可渲染的新算子模板如add_custom.yamlreferences/适配规则手册sk-adaptation-cookbook.md。所有命令统一通过一个入口脚本调用python3 skills_root/sk-operator-codegen/scripts/operator_codegen.py subcommand ...其中skills_root指包含sk-operator-codegen目录的 skill 根路径。在仓库中对应 operator_codegen.py其main()通过build_parser()构建子命令分发见源码末尾的main()与__main__入口。二、源码形态识别detect-sk-form适配前首先要回答一个问题这份算子源码到底是什么形态detect-sk-form子命令负责把源码分类并输出operator-sk-form-analysis.json。python3 skills_root/sk-operator-codegen/scripts/operator_codegen.py detect-sk-form \ $OPERATOR_ASSET \ --output-dir build/examples/sk-codegen/detect支持四类形态源码实现见 operator_codegen.py 的cmd_detect_sk_form逐文件调用sk_codegen_lib.detect_sk_form后汇总输入形态典型特征处理结果none普通__global__kernel无 SK 标记生成 Args struct、模板化__sk__和SK_BINDcurrent-sk-bind已包含当前框架需要的 bind 结构__sk__、SK_BIND、CommArgs/SkSystemArgs等按字节复制源码标记already_currentpartial部分 SK 特征存在但不完整不猜测输出codegen.unknown-sk-form等人工处理项unknown缺少入口、规格或语义信息输出诊断结果交由人工处理当多个源文件形态不一致时总览结果会被归类为mixed。detect_sk_form的标记检测逻辑见 sk_codegen_lib.py会同时检查__sk__关键字、SK_BIND宏以及 SK 参数结构体特征CommArgs、SkSystemArgs、__gm__ uint64_t *param并给出每项标记的置信度与证据文件。三、核心命令adapt-sk-from-global 自动生成 SK bind形态识别确认算子属于none普通__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.json该命令统一处理已分类的输入形态普通或可修复的非 SK 源码先走 codegen 拥有的预适配自动修复如自动补#include kernel_operator.h、清理仅__global__可用的 kernel task 宏、为TPipe补充Destroy等见 sk_codegen_lib.py 的adapt_source_text再生成当前 SK bind。模板化 global 会保留模板参数并追加splitidx只有字段类型依赖模板参数时Args struct 才模板化。3.1 生成包遵循 aclgraph-canonical 布局每个算子适配后输出一棵标准的 aclgraph-canonical 目录树operator-sk-adapted/ csrc/op.asc # 适配后的算子源码含 SK bind csrc/pybind11.asc # pybind11 绑定源码 op_extension/__init__.py op_extension/_torch_library.py setup.py # wheel 构建脚本3.2 IO contract多 tensor 参数的语义契约当 kernel 有多个 tensor-like 参数时--io-contract FILE是必选项。固化脚本不会根据变量名或参数顺序猜测输入输出必须由用户、adapter skill 或上游资产契约明确说明每个 tensor-like 参数属于inputs、outputs还是workspaces以及 pybind 单返回值应该返回哪个 tensor。最小格式如下{ schema_version: 1, entries: { add_custom: { inputs: [x, y], outputs: [z], pybind_return_tensor: z } } }契约解析逻辑在 operator_codegen.py 的_read_pybind_io_contract中实现支持以下完整字段inputs/outputs/workspacestensor-like 参数的归属声明workspaces用于GM_ADDR workspace、GM_ADDR tiling这类地址参数注意这类地址参数应放入workspaces不要声明成 host structpybind_return_tensorpybind 单返回值对应的输出 tensor必须在outputs中列出parameters.name.kind运行时参数种类支持tensor、tensor_list、scalar、host_struct四种struct-valued 运行时参数例如tiling: {kind: host_struct}必须显式声明否则契约解析会直接报错parameters.name.nullable布尔值标记参数是否可空launch.block_dim正整数声明算子 launch 的 block 维度compile.defines/compile.options编译宏与编译选项宏会自动补-D前缀runtime_wrapperTensorList 等复杂运行时参数的自定义 wrapper 声明见下文「图捕获前准备状态」。3.3 契约缺失或错误时的明确结论契约驱动意味着「宁缺毋滥」缺失或不完整时不会做 best-effort 兜底而是输出needs-human或具体 codegen finding多 tensor-like 参数但无--io-contract→ 返回needs-human在operator-sk-adapted.json中给出codegen.pybind-return-tensor-unresolved契约匹配但遗漏了某个 tensor-like 参数 →codegen.io-contract-tensor-incomplete存在 host-struct 参数但未声明 →codegen.runtime-parameter-contract-required契约声明了未知参数、参数 kind 与源码检测不一致、bind_target不匹配、TensorList 参数缺少 runtime wrapper → 分别输出codegen.io-contract-parameter-invalid、codegen.io-contract-parameter-kind-mismatch、codegen.io-contract-bind-target-mismatch、codegen.tensor-list-runtime-wrapper-required。这些 preflight 检查在cmd_adapt_sk_from_global中通过apply_io_contract逐 entry 执行见 operator_codegen.py。四、SK bind 生成形态详解对每一种支持的源码形态adapter 要么生成当前 SK binding要么返回明确的人工处理项。对于干净的none形态会在原始__global__函数之后按顺序生成三样东西规则契约见 sk-adaptation-cookbook.md代码实现见 sk_codegen_lib.py 的render_sk_adaptation。4.1 Args structkernel 参数的统一载体struct NameCamelArgs { c_type param; // 每个 kernel 参数对应一个字段保持原始顺序 // ... };命名为NameCamelArgskernel 名转 CamelCase 后加Args后缀每个 kernel 参数对应一个字段保持原始顺序C 类型小于 4 字节的字段int8_t、uint8_t、int16_t、uint16_t、bool自动加alignas(4)保证结构体对齐语义源码中_is_small_int_type判断 align_prefix拼接若原始 kernel 是模板函数只有字段类型依赖模板参数时 Args struct 才模板化只影响 body 或 kernel type 的模板参数保留在 SK 函数上不改变 runtime 参数包布局。4.2 模板化__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 默认不改动逐行解包args后按原样保留只有 body 引用了AscendC::GetBlockNum()时才注入sysArgs参数并把调用改写为sysArgs-skNumBlocks见 sk_codegen_lib.py 的正则改写逻辑模板参数追加uint32_t splitidx用于 SK 的多核拆分split标识。4.3 SK_BIND 注册语句SK_BIND(orig, mask, name_sk0, name_sk1, name_sk2, name_sk3)mask默认是4DCCI/早启能力位允许值0..70表示没有能力 bit1/2/4是 bit flag源码中校验mask in range(0, 8)--num-splits控制绑定多少个name_skN符号范围1..4源码中校验1 num_splits 4超出即抛错原始__global__函数必须保持不变。4.4 Kernel 类型映射__sk__函数会继承原始 kernel 的执行类型 qualifier映射规则如下源码见map_kernel_type_for_sk原始 qualifierSK qualifier__vector____vector____cube____cube____mix__(c, v)general__mix__(c, v)__mix__(1, 0)__cube__特殊情况__mix__(0, 1)__vector__特殊情况bare__aicore____aicore__4.5 sysArgs 注入策略--with-sys-argsauto是默认值只有原始 body 包含AscendC::GetBlockNum()时才注入sk::SkSystemArgs *sysArgs--with-sys-argsalways/never可以强制选择实现见 sk_codegen_lib.py 的sys_args_mode三分支。注入后使用当前 API 名sysArgs-skNumBlocks/sysArgs-SkGetNumBlocks()。历史写法skBlockNum/SkGetBlockNum在当前 CANN 头文件下会编译失败sk-operator-validate --rule-pack spec会将其标记为sk.sys-args-api-current并支持自动重命名修复。五、图捕获前准备状态codegen 的运行时状态规则部分算子需要在图捕获前准备运行状态例如 TensorList descriptor、workspace tail 元数据或持久缓存。本 skill 对这类场景有明确的硬性规则见 SKILL.md 的「运行时状态规则」一节tensor-like 参数角色不明确时必须要求 IO contractcapture 前必须准备的 runtime state例如 TensorList descriptor 或持久 workspace 元数据必须由显式契约表达并由 runtime wrapper 消费生成 wrapper不得在 captured forward 路径里插入隐藏 helper kernel来临时制造 descriptor 或状态descriptor 顺序必须由 contract 驱动并与声明的 TensorList 参数校验一致不一致时是生成错误不做 best-effort fallback。在 IO contract 中TensorList 参数需要声明kind tensor_list同时通过runtime_wrapper声明自定义 wrapper 源码与入口例如{ schema_version: 1, entries: { concat_custom: { inputs: [inputs_list], outputs: [out], pybind_return_tensor: out, parameters: { inputs_list: { kind: tensor_list } }, runtime_wrapper: { source: runtime_wrapper.cpp, entry: concat_wrapper, tensor_list_descriptor_strategy: prepared_workspace_tail, prepare_entry: concat_prepare, descriptor_bytes: 64, descriptor_order: [inputs_list] } } } }tensor_list_descriptor_strategy目前仅支持prepared_workspace_tail一种取值且一旦声明就必须同时提供prepare_entry、正整数descriptor_bytes与descriptor_order校验见 operator_codegen.py 的_runtime_wrapper_contract。wrapper 的source必须是 asset 内部的相对路径禁止绝对路径与..越界entry必须是合法 C 标识符。这条规则适用于所有需要 prepared runtime state 的算子不是某个样例的特殊逻辑。六、多算子聚合aggregate-sk-adapted单个算子适配后仍然各自独立如果希望把它们打包进一个wheel需要聚合步骤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 布局operator-sk-adapted/csrc/*.asc、csrc/pybind11.asc、op_extension/、setup.py但有以下变化pybind11.asc和_torch_library.py会注册所有算子的 entry聚合内 entry 名必须唯一重复会作为错误处理生成的 pybind 层为每个算子暴露面向用户的 bind target entryrun_op通过torch.library注册differential validation 在 baseline 和 SK context 下复用同一个入口聚合setup.py保持 Python import 包名为op_extension同时用用户指定的 distribution name--aggregate-wheel-name和 version--package-version生成 wheel 文件名wheel 的 native module 仍按entry x arch拆分详见下文第八节。七、生成 standalone compare 工程聚合完成后可以生成一个独立的对比验证工程用于在真实设备上对比 baseline 与 SK 的执行结果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可直接编译的 CMake 工程复制后的 adapted csrcoperator-sk-standalone-verify.json工程描述。生成的源码会输出launch_op_baseline和launch_op_sk两个 wrapper两者调用同一个bind_target由 runtime context 区分 baseline 与 SK保证对比口径一致。7.1 设备与 arch 的边界约束如果 fixture 声明了 device buffers/scalars会分配独立的baseline/SK buffer、回拷可比输出并输出byte/hash 对比结果没有显式 device plan 时真实设备运行返回skipped-insufficient-runtime-spec不能伪造通过standalone 的--npu-arch必须显式或只能由--target-chip解析到唯一的 source-backed arch否则输出needs-target-arch不会静默回退到某个默认 arch。八、多 arch 支持SK_NPU_ARCHS 与运行时 arch 选择生成的 ACLGraph wheel 支持多芯片版本原生产物arch 处理贯穿构建期与运行期构建期优先读取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 编译数自动映射SK_NPU_ARCHS与SK_BISHENG_JOBS都未设置时只使用有官方源码依据的当前环境检测。目前自动映射只覆盖Ascend950*/ torch_npu SoC enum260到dav-3510其他芯片不会静默 fallback需显式设置SK_NPU_ARCHS映射实现见 sk_codegen_lib.py 的_detect_soc_version/map_soc_version_to_arch运行期优先读取SK_ACLGRAPH_NPU_ARCH未设置时才尝试有来源依据的 SoC 自动映射。即使 wheel 里只有一个.so也不会在无法确认目标架构时静默选择。九、自动修复apply-remediation静态检查如sk-operator-validate产出的 findings 中有一部分是机器可自动修复的。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.json支持的自动修复种类对应sk_codegen_lib.py中的AUTO_REMEDIATION_KINDS常量rename-symbol符号重命名remove-line-containing删除包含指定内容的行add-include添加头文件包含replace-pattern正则模式替换。不可自动修复的项会作为人工处理项输出交由开发者处理。新增自动修复规则的方法是扩展scripts/sk_codegen_lib.py中的AUTO_REMEDIATION_KINDS并在apply_remediation中增加对应处理分支。十、模板生成从干净输入开始闭环流水线需要一个「干净输入」起点generate-from-template和list-templates提供这一能力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-out模板渲染templates/id.yaml。仓库内置的add_custom模板见 add_custom.yaml是一个最小化的非 SK elementwise add 算子生成干净的__global__ __vector__kernel可以直接喂给adapt-sk-from-global派生 SK 适配id: add_custom parameters: - name: dtype kind: choice default: float16 allowed: [float16, float32, int32] files: - path: csrc/$id.asc template: | #include kernel_operator.h // ... class KernelAdd 与 extern C __global__ __vector__ void $id(...) ...模板使用string.Template风格占位符$id、$dtype、$ctype因此 kernel 代码中的花括号无需转义。参数支持choice类型并声明allowed取值集合。新增基础算子模板的方法是新增templates/id.yaml。十一、输出约定与流水线落位生成阶段通常产出四类文件适配后的源码目录描述算子、输入、输出、构建配置的manifest聚合目录_aggregate供 pybind、wheel 和 standalone 阶段继续使用诊断 JSON用于说明为什么某个算子不能自动生成。在总流水线中这些文件会落到约定目录inputs/outputs分别承接上游输入与下游产物01-detect-sk-form/op/{inputs,outputs} 02-adapt-sk-from-global/op/{inputs,outputs} 02-adapt-sk-from-global/_aggregate/{inputs,outputs}十二、交付边界与其他 skill 的衔接本 skill 的输出交给下游三个 skill 继续消费见 SKILL.md 的「交付边界」一节sk-operator-validate执行 contract/spec/compat 规则包输出统一 findingssk-operator-build-package消费operator-sk-adapted.json和operator-sk-adapted/生成 pybind binding、wheel并构建 standalone compare 工程sk-operator-pipeline run-sk-pipeline编排完整闭环。需要说明的是端到端场景优先使用sk-operator-pipeline run-sk-pipeline本工具更适合单独定位生成阶段的问题见 README.md。何时直接运行本工具只想确认一个算子是否能被识别生成阶段失败需要单独重跑adapt-sk-from-global想检查聚合后的目录是否满足后续 pybind/wheel 构建要求需要用apply-remediation对静态检查结果做自动修复尝试。十三、历史 scaffold 入口兼容保留intake、plan、analyze-sk-conversion、adapt-sk-binding-scaffold、generate-sk-source-scaffold作为历史 scaffold 入口仍保留详细说明见references/workflow.md。新代码路径统一收敛到上文所述的各子命令。十四、验证与快速自检运行入口脚本的--help可以快速验证 skill 环境是否可用python3 skills_root/sk-operator-codegen/scripts/operator_codegen.py --help一条完整的最小验证路径是# 1. 用内置模板生成干净输入 python3 skills_root/sk-operator-codegen/scripts/operator_codegen.py generate-from-template \ add_custom --param dtypefloat16 \ --output-dir build/examples/sk-codegen/template-out # 2. 识别形态预期none python3 skills_root/sk-operator-codegen/scripts/operator_codegen.py detect-sk-form \ build/examples/sk-codegen/template-out --output-dir build/examples/sk-codegen/detect # 3. 生成 SK bind单 tensor 输出可不带 --io-contract python3 skills_root/sk-operator-codegen/scripts/operator_codegen.py adapt-sk-from-global \ build/examples/sk-codegen/template-out --output-dir build/examples/sk-codegen/adapted/add_custom总结sk-operator-codegen以「契约驱动、绝不猜测」为设计主线形态识别负责分类IO contract 负责把 tensor 语义与运行时状态显式化adapt-sk-from-global生成标准的 Args struct __sk__SK_BIND三件套aggregate-sk-adapted与generate-standalone-compare打通多算子打包与真实设备对比验证而模板与AUTO_REMEDIATION_KINDS则提供了可持续扩展的入口。对于希望把普通 AscendC 算子接入 SuperKernel 的开发者这条从detect-sk-form到generate-standalone-compare的链路就是最直接的实操路径对于框架开发者sk-adaptation-cookbook.md 中关于 Args struct 对齐、kernel 类型映射、sysArgs 注入的规则则是理解生成产物正确性的权威依据。【免费下载链接】graph-autofusionGraph-autofusion 是一个面向昇腾Ascend芯片的轻量级、解耦式组件集合旨在通过自动融合技术加速模型执行。 目前已开源 SuperKernel 组件和 Autofuse 组件未来将持续开放更多自动融合相关模块。项目地址: https://gitcode.com/cann/graph-autofusion创作声明：本文部分内容由AI辅助生成（AIGC），仅供参考