尧图网页设计 
CANN graph-autofusion SuperKernel 算子 SK 适配规则手册从__global__kernel 到 SK_BIND 绑定【免费下载链接】graph-autofusionGraph-autofusion 是一个面向昇腾Ascend芯片的轻量级、解耦式组件集合旨在通过自动融合技术加速模型执行。 目前已开源 SuperKernel 组件和 Autofuse 组件未来将持续开放更多自动融合相关模块。项目地址: https://gitcode.com/cann/graph-autofusion导读本文是 CANN graph-autofusion 开源仓库中sk-operator-codegenskill 的适配规则手册adapt-sk-from-global编码规则的快速参考系统讲解如何把普通 AscendC__global__kernel 自动适配为 SuperKernelSKbinding 形态从 Args struct、模板化__sk__函数到SK_BIND语句的生成规则并覆盖多算子聚合渲染、输入形态分类、kernel 类型映射与sysArgs注入时机。读完本文你将掌握 SK binding 的完整代码形态契约、operator_codegen.py各子命令的调用方式以及源码层面对每条规则的实现依据可直接用于日常算子 SK 化适配与问题定位。一、适配产物一次__global__到 SK binding 的完整转换对每一种支持的源码形态adapt-sk-from-globaladapter 要么生成当前 SK binding要么返回明确的人工处理项。对于干净的非 SK__global__kernel会在原始函数之后按顺序生成三块内容原始__global__函数本身必须保持不变。这三块内容在源码中的拼接位置与顺序由 sk_codegen_lib.py 中的adapt_source_text负责它遍历每个__global__入口找到函数结尾后依次插入// ---- SK adaptation (auto-generated) ----标记块。1. Args structkernel 参数的结构化封装命名为NameCamelArgs每个 kernel 参数对应一个字段保持原始顺序。示例对应add_custom这类 elementwise 算子struct AddCustomArgs { GM_ADDR x; GM_ADDR y; GM_ADDR z; uint32_t totalLength; };关键规则C 类型小于 4 字节的字段例如int8_t、uint8_t、int16_t、uint16_t、bool必须使用alignas(4)前缀。这一规则对应 sk_codegen_lib.py 中的SMALL_INT_TYPES常量集合渲染时通过_is_small_int_type(p.c_type)判断并加上alignas(4)前缀见 sk_codegen_lib.py。这是 ABI 层面的硬性要求用于保证 runtime 参数包布局的稳定性。模板函数处理只有字段类型依赖模板参数时Args struct 才模板化只影响 body 或 kernel type 的模板参数保留在 SK 函数上不改变 runtime 参数包布局。渲染逻辑见 sk_codegen_lib.py通过_args_template_params_for_fields仅筛选出影响字段类型的模板参数。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 默认不改动verbatim 复制只做两处例外改写一是 body 引用了AscendC::GetBlockNum()时注入sysArgs参数并把调用改写为sysArgs-skNumBlocks正则替换实现见 sk_codegen_lib.py二是对仅适用于__global__的 kernel task type 宏做剥离、对TPipe生命周期补destroy调用见_strip_global_only_kernel_task_type_macros与_ensure_tpipe_destroy_without_pipe_all。每个参数生成一行解包语句c_type param args-param;保持与原始 kernel 相同的参数名与类型。模板参数uint32_t splitidx恒被追加与原始模板参数如有一起构成最终模板形参列表渲染见 sk_codegen_lib.py。3.SK_BIND语句SK_BIND(orig, mask, name_sk0, name_sk1, name_sk2, name_sk3)SK_BIND将原始算子在 SK runtime 侧的 bind target 与若干 split 符号绑定mask默认是4DCCI即 Dump/Compare/Check/Inspect 类能力 bit。允许值是0..7其中0表示没有能力 bit1/2/4是 bit flag可组合。代码层面的取值范围校验在 sk_codegen_lib.pymask must be in 0..7。--num-splits控制绑定多少个name_skN符号范围1..4校验在 sk_codegen_lib.py。CLI 默认值为4见 operator_codegen.py。生成的宏文本形如SK_BIND(bind_symbol, mask, split_args);其中split_args为name_sk0, name_sk1, ...的逗号拼接见 sk_codegen_lib.py。二、多算子聚合渲染从单算子树到聚合 wheel单算子适配产物aclgraph-canonical 布局单算子适配仍然为每个 asset 写一个 aclgraph-canonical treeoperator-sk-adapted/ csrc/op.asc csrc/pybind11.asc op_extension/__init__.py op_extension/_torch_library.py setup.py聚合树与 entry 唯一性aggregate-sk-adapted消费多个这样的输出渲染一个聚合 tree包含所有csrc/op.asc、一个pybind11.asc、一个_torch_library.py和一个setup.py。约束聚合内 entry 名必须唯一否则 pybind 层无法为每个算子生成可区分的 bind target。生成的 pybind 层为每个算子暴露面向用户的 bind target entry函数作用run_op通过torch.library注册的 SK-facing bind targetdifferential validation 在 baseline 和 SK context 下复用同一个入口聚合setup.py保持 Python import 包名为op_extension同时用用户指定的 distribution name 和 version 生成 wheel 文件名对应--aggregate-wheel-name与--package-version参数。聚合命令示例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多 arch 原生产物支持生成的 ACLGraph wheel 包支持多芯片版本原生产物构建时优先读取SK_NPU_ARCHS可用逗号或分号传多个值例如SK_NPU_ARCHSdav-2201,dav-3510。wheel 内 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也不会在无法确认目标架构时静默选择。三、输入形态分类与适配行为detect-sk-form子命令将算子源码分类为四种形态并输出operator-sk-form-analysis.json输入形态适配行为none生成 Args struct、模板化__sk__和SK_BIND。可修复none在临时副本上执行 codegen 拥有的预适配自动修复再生成当前 SK binding。current-sk-bind按字节复制源码并标记为already_current。partial/unknown不猜测输出codegen.unknown-sk-form等人工处理项。形态判定的源码依据在 sk_codegen_lib.pycurrent-sk-bind要求同时存在__sk__与SK_BIND模板风格partial指混合信号例如有__sk__但没有SK_BIND。operator_codegen.py中的_collect_sk_markersoperator_codegen.py通过扫描__sk__、SK_BIND、CommArgs/SkSystemArgs/__gm__ uint64_t *param标记来辅助判定。对于不可自动处理的输入adapter 会输出明确的诊断而不是猜测这是本 skill 的本地行为契约核心宁可交给人工也不生成可能错误的绑定。四、Kernel 类型映射规则原始__global__函数的 kernel-type qualifier 在生成__sk__函数时按下表映射实现见 sk_codegen_lib.py 的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__映射规则要点__mix__特殊情形优先当__mix__(c, v)中c 1 v 0时归一化为__cube__当c 0 v 1时归一化为__vector__其余__mix__(c, v)保持原样。无法从 qualifier 推导出任何已知类型时函数抛出ValueErrorcannot map kernel type from qualifiers避免生成非法 SK 签名。五、sysArgs注入时机与 API 命名约束三种模式模式行为--with-sys-argsauto默认只有原始 body 包含AscendC::GetBlockNum()时才注入sk::SkSystemArgs *sysArgs。--with-sys-argsalways无条件注入。--with-sys-argsnever强制不注入。实现上sys_args_mode的三种取值判定在 sk_codegen_lib.pyauto模式直接读取解析阶段得到的entry.uses_get_block_num标志。注入后的 API 命名注入后必须使用当前 API 名sysArgs-skNumBlockssysArgs-SkGetNumBlocks()历史命名不可用skBlockNum/SkGetBlockNum在当前 CANN 头文件下会编译失败。sk-operator-validate --rule-pack spec会将其标记为sk.sys-args-api-current并支持自动重命名修复——对应的新旧名称映射对在 sk_codegen_lib.py 中定义(skBlockNum, skNumBlocks)与(SkGetBlockNum, SkGetNumBlocks)。六、完整命令行流程与验证端到端常用命令识别单个算子形态python3 skills_root/sk-operator-codegen/scripts/operator_codegen.py detect-sk-form \ $OPERATOR_ASSET \ --output-dir build/examples/sk-codegen/detect把普通__global__算子适配为 SK bindpython3 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生成 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基于检查结果应用自动修复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查看模板能力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-outstandalone compare 的架构约束standalone compare 必须有明确的 NPU arch优先显式传--npu-arch未传时只会在--target-chip能通过官方来源映射到唯一 arch 时继续生成可编译 CMake否则输出needs-target-arch不会静默回退到某个默认 arch。若 fixture 声明 device buffers/scalars会分配独立 baseline/SK buffer、回拷可比输出并输出 byte/hash 对比结果没有显式 device plan 时真实设备运行返回skipped-insufficient-runtime-spec不伪造通过。输出约定与流水线落位生成阶段通常产出适配后的源码目录、描述算子/输入/输出/构建配置的 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}七、IO 契约为什么不猜变量名--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 } } }契约规则要点解析与校验实现在 operator_codegen.py如果一个 entry 有多个 tensor-like 参数但没有匹配的--io-contractadapt-sk-from-global返回needs-human并在operator-sk-adapted.json中给出codegen.pybind-return-tensor-unresolved。如果契约匹配但遗漏了某个 tensor-like 参数给出codegen.io-contract-tensor-incomplete。struct-valued 运行时参数需要在parameters中声明例如tiling: {kind: host_struct}GM_ADDR workspace或GM_ADDR tiling这类地址参数仍应放入workspaces不要声明成 host struct。合法kind取值包括tensor、tensor_list、scalar、host_struct校验见 operator_codegen.pyparameters还支持nullable布尔声明。图捕获前准备状态runtime state规则有些算子需要在图捕获前准备运行状态例如 TensorList descriptor、workspace tail 元数据或持久缓存。生成规则是这些状态必须来自显式 contract不从源码变量名或参数顺序推断。runtime wrapper 只能消费已经准备好的状态。不允许在 forward/capture 路径中生成隐藏 helper kernel 来临时准备状态。descriptor 顺序必须和 contract 中声明的 TensorList 参数顺序一致不一致时返回needs-human或生成错误。这条规则适用于所有需要 prepared runtime state 的算子不是某个样例的特殊逻辑。契约中对应字段为runtime_wrapper含source、entry、tensor_list_descriptor_strategy: prepared_workspace_tail、prepare_entry、descriptor_bytes、descriptor_order解析实现见 operator_codegen.py。八、自动修复与模板扩展点自动修复项apply-remediation对静态检查 findings 应用机器可修复项支持四种 kind定义见 sk_codegen_lib.pyrename-symbol符号重命名如历史 sysArgs API 名 → 当前 API 名。remove-line-containing删除包含指定内容的行。add-include补充缺失头文件。replace-pattern模式替换。不可自动修复项会作为人工处理项输出。新自动修复规则的扩展方式是向AUTO_REMEDIATION_KINDS增加 kind并在apply_remediation中增加处理分支。模板扩展点新基础算子模板新增templates/id.yaml。仓库自带的 add_custom.yaml 是一个最小非 SK elementwise add 算子模板渲染出一个干净的__global__ __vector__kernel含dtype参数可选float16/float32/int32可直接作为adapt-sk-from-global的干净输入演示闭环。新自动修复规则扩展scripts/sk_codegen_lib.py中的AUTO_REMEDIATION_KINDS。何时直接运行本工具只想确认一个算子是否能被识别detect-sk-form。生成阶段失败需要单独重跑adapt-sk-from-global。想检查聚合后的目录是否满足后续 pybind/wheel 构建要求。需要用apply-remediation对静态检查结果做自动修复尝试。验证命令python3 skills_root/sk-operator-codegen/scripts/operator_codegen.py --help九、在 skill 流水线中的位置与交付边界本 skill 的输出交给下游三个 skill 消费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 run-sk-pipeline编排完整闭环。端到端场景优先使用sk-operator-pipeline run-sk-pipelinesk-operator-codegen适合单独定位生成阶段问题。完整命令与产物说明见 SKILL.md 与 README.md。此外intake、plan、analyze-sk-conversion、adapt-sk-binding-scaffold、generate-sk-source-scaffold等历史 scaffold 入口仍保留供兼容场景使用。十、总结本手册作为本地行为契约sk-adaptation-cookbook.md之所以被定义为本地行为契约是因为它把 adapter 的行为约束到可测试、可复现的程度能自动生成的形态严格生成、不能确定的形态明确上报同时以原始__global__函数保持不变不猜变量名不静默回退 arch等硬规则保护生成产物的正确性。完整规格和边界情况以脚本实现、测试和本手册三者共同约束其中 sk_codegen_lib.py 是适配渲染器Args struct __sk__template SK_BIND与自动修复器的实现主体operator_codegen.py 是 CLI 入口与 IO 契约解析器二者共同构成本文所有规则的可执行依据。【免费下载链接】graph-autofusionGraph-autofusion 是一个面向昇腾Ascend芯片的轻量级、解耦式组件集合旨在通过自动融合技术加速模型执行。 目前已开源 SuperKernel 组件和 Autofuse 组件未来将持续开放更多自动融合相关模块。项目地址: https://gitcode.com/cann/graph-autofusion创作声明：本文部分内容由AI辅助生成（AIGC），仅供参考 CANN graph-autofusion SuperKernel 算子 SK 适配规则手册：从 `__global__` kernel 到 SK_BIND 绑定
CANN graph-autofusion SuperKernel 算子 SK 适配规则手册从__global__kernel 到 SK_BIND 绑定【免费下载链接】graph-autofusionGraph-autofusion 是一个面向昇腾Ascend芯片的轻量级、解耦式组件集合旨在通过自动融合技术加速模型执行。 目前已开源 SuperKernel 组件和 Autofuse 组件未来将持续开放更多自动融合相关模块。项目地址: https://gitcode.com/cann/graph-autofusion导读本文是 CANN graph-autofusion 开源仓库中sk-operator-codegenskill 的适配规则手册adapt-sk-from-global编码规则的快速参考系统讲解如何把普通 AscendC__global__kernel 自动适配为 SuperKernelSKbinding 形态从 Args struct、模板化__sk__函数到SK_BIND语句的生成规则并覆盖多算子聚合渲染、输入形态分类、kernel 类型映射与sysArgs注入时机。读完本文你将掌握 SK binding 的完整代码形态契约、operator_codegen.py各子命令的调用方式以及源码层面对每条规则的实现依据可直接用于日常算子 SK 化适配与问题定位。一、适配产物一次__global__到 SK binding 的完整转换对每一种支持的源码形态adapt-sk-from-globaladapter 要么生成当前 SK binding要么返回明确的人工处理项。对于干净的非 SK__global__kernel会在原始函数之后按顺序生成三块内容原始__global__函数本身必须保持不变。这三块内容在源码中的拼接位置与顺序由 sk_codegen_lib.py 中的adapt_source_text负责它遍历每个__global__入口找到函数结尾后依次插入// ---- SK adaptation (auto-generated) ----标记块。1. Args structkernel 参数的结构化封装命名为NameCamelArgs每个 kernel 参数对应一个字段保持原始顺序。示例对应add_custom这类 elementwise 算子struct AddCustomArgs { GM_ADDR x; GM_ADDR y; GM_ADDR z; uint32_t totalLength; };关键规则C 类型小于 4 字节的字段例如int8_t、uint8_t、int16_t、uint16_t、bool必须使用alignas(4)前缀。这一规则对应 sk_codegen_lib.py 中的SMALL_INT_TYPES常量集合渲染时通过_is_small_int_type(p.c_type)判断并加上alignas(4)前缀见 sk_codegen_lib.py。这是 ABI 层面的硬性要求用于保证 runtime 参数包布局的稳定性。模板函数处理只有字段类型依赖模板参数时Args struct 才模板化只影响 body 或 kernel type 的模板参数保留在 SK 函数上不改变 runtime 参数包布局。渲染逻辑见 sk_codegen_lib.py通过_args_template_params_for_fields仅筛选出影响字段类型的模板参数。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 默认不改动verbatim 复制只做两处例外改写一是 body 引用了AscendC::GetBlockNum()时注入sysArgs参数并把调用改写为sysArgs-skNumBlocks正则替换实现见 sk_codegen_lib.py二是对仅适用于__global__的 kernel task type 宏做剥离、对TPipe生命周期补destroy调用见_strip_global_only_kernel_task_type_macros与_ensure_tpipe_destroy_without_pipe_all。每个参数生成一行解包语句c_type param args-param;保持与原始 kernel 相同的参数名与类型。模板参数uint32_t splitidx恒被追加与原始模板参数如有一起构成最终模板形参列表渲染见 sk_codegen_lib.py。3.SK_BIND语句SK_BIND(orig, mask, name_sk0, name_sk1, name_sk2, name_sk3)SK_BIND将原始算子在 SK runtime 侧的 bind target 与若干 split 符号绑定mask默认是4DCCI即 Dump/Compare/Check/Inspect 类能力 bit。允许值是0..7其中0表示没有能力 bit1/2/4是 bit flag可组合。代码层面的取值范围校验在 sk_codegen_lib.pymask must be in 0..7。--num-splits控制绑定多少个name_skN符号范围1..4校验在 sk_codegen_lib.py。CLI 默认值为4见 operator_codegen.py。生成的宏文本形如SK_BIND(bind_symbol, mask, split_args);其中split_args为name_sk0, name_sk1, ...的逗号拼接见 sk_codegen_lib.py。二、多算子聚合渲染从单算子树到聚合 wheel单算子适配产物aclgraph-canonical 布局单算子适配仍然为每个 asset 写一个 aclgraph-canonical treeoperator-sk-adapted/ csrc/op.asc csrc/pybind11.asc op_extension/__init__.py op_extension/_torch_library.py setup.py聚合树与 entry 唯一性aggregate-sk-adapted消费多个这样的输出渲染一个聚合 tree包含所有csrc/op.asc、一个pybind11.asc、一个_torch_library.py和一个setup.py。约束聚合内 entry 名必须唯一否则 pybind 层无法为每个算子生成可区分的 bind target。生成的 pybind 层为每个算子暴露面向用户的 bind target entry函数作用run_op通过torch.library注册的 SK-facing bind targetdifferential validation 在 baseline 和 SK context 下复用同一个入口聚合setup.py保持 Python import 包名为op_extension同时用用户指定的 distribution name 和 version 生成 wheel 文件名对应--aggregate-wheel-name与--package-version参数。聚合命令示例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多 arch 原生产物支持生成的 ACLGraph wheel 包支持多芯片版本原生产物构建时优先读取SK_NPU_ARCHS可用逗号或分号传多个值例如SK_NPU_ARCHSdav-2201,dav-3510。wheel 内 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也不会在无法确认目标架构时静默选择。三、输入形态分类与适配行为detect-sk-form子命令将算子源码分类为四种形态并输出operator-sk-form-analysis.json输入形态适配行为none生成 Args struct、模板化__sk__和SK_BIND。可修复none在临时副本上执行 codegen 拥有的预适配自动修复再生成当前 SK binding。current-sk-bind按字节复制源码并标记为already_current。partial/unknown不猜测输出codegen.unknown-sk-form等人工处理项。形态判定的源码依据在 sk_codegen_lib.pycurrent-sk-bind要求同时存在__sk__与SK_BIND模板风格partial指混合信号例如有__sk__但没有SK_BIND。operator_codegen.py中的_collect_sk_markersoperator_codegen.py通过扫描__sk__、SK_BIND、CommArgs/SkSystemArgs/__gm__ uint64_t *param标记来辅助判定。对于不可自动处理的输入adapter 会输出明确的诊断而不是猜测这是本 skill 的本地行为契约核心宁可交给人工也不生成可能错误的绑定。四、Kernel 类型映射规则原始__global__函数的 kernel-type qualifier 在生成__sk__函数时按下表映射实现见 sk_codegen_lib.py 的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__映射规则要点__mix__特殊情形优先当__mix__(c, v)中c 1 v 0时归一化为__cube__当c 0 v 1时归一化为__vector__其余__mix__(c, v)保持原样。无法从 qualifier 推导出任何已知类型时函数抛出ValueErrorcannot map kernel type from qualifiers避免生成非法 SK 签名。五、sysArgs注入时机与 API 命名约束三种模式模式行为--with-sys-argsauto默认只有原始 body 包含AscendC::GetBlockNum()时才注入sk::SkSystemArgs *sysArgs。--with-sys-argsalways无条件注入。--with-sys-argsnever强制不注入。实现上sys_args_mode的三种取值判定在 sk_codegen_lib.pyauto模式直接读取解析阶段得到的entry.uses_get_block_num标志。注入后的 API 命名注入后必须使用当前 API 名sysArgs-skNumBlockssysArgs-SkGetNumBlocks()历史命名不可用skBlockNum/SkGetBlockNum在当前 CANN 头文件下会编译失败。sk-operator-validate --rule-pack spec会将其标记为sk.sys-args-api-current并支持自动重命名修复——对应的新旧名称映射对在 sk_codegen_lib.py 中定义(skBlockNum, skNumBlocks)与(SkGetBlockNum, SkGetNumBlocks)。六、完整命令行流程与验证端到端常用命令识别单个算子形态python3 skills_root/sk-operator-codegen/scripts/operator_codegen.py detect-sk-form \ $OPERATOR_ASSET \ --output-dir build/examples/sk-codegen/detect把普通__global__算子适配为 SK bindpython3 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生成 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基于检查结果应用自动修复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查看模板能力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-outstandalone compare 的架构约束standalone compare 必须有明确的 NPU arch优先显式传--npu-arch未传时只会在--target-chip能通过官方来源映射到唯一 arch 时继续生成可编译 CMake否则输出needs-target-arch不会静默回退到某个默认 arch。若 fixture 声明 device buffers/scalars会分配独立 baseline/SK buffer、回拷可比输出并输出 byte/hash 对比结果没有显式 device plan 时真实设备运行返回skipped-insufficient-runtime-spec不伪造通过。输出约定与流水线落位生成阶段通常产出适配后的源码目录、描述算子/输入/输出/构建配置的 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}七、IO 契约为什么不猜变量名--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 } } }契约规则要点解析与校验实现在 operator_codegen.py如果一个 entry 有多个 tensor-like 参数但没有匹配的--io-contractadapt-sk-from-global返回needs-human并在operator-sk-adapted.json中给出codegen.pybind-return-tensor-unresolved。如果契约匹配但遗漏了某个 tensor-like 参数给出codegen.io-contract-tensor-incomplete。struct-valued 运行时参数需要在parameters中声明例如tiling: {kind: host_struct}GM_ADDR workspace或GM_ADDR tiling这类地址参数仍应放入workspaces不要声明成 host struct。合法kind取值包括tensor、tensor_list、scalar、host_struct校验见 operator_codegen.pyparameters还支持nullable布尔声明。图捕获前准备状态runtime state规则有些算子需要在图捕获前准备运行状态例如 TensorList descriptor、workspace tail 元数据或持久缓存。生成规则是这些状态必须来自显式 contract不从源码变量名或参数顺序推断。runtime wrapper 只能消费已经准备好的状态。不允许在 forward/capture 路径中生成隐藏 helper kernel 来临时准备状态。descriptor 顺序必须和 contract 中声明的 TensorList 参数顺序一致不一致时返回needs-human或生成错误。这条规则适用于所有需要 prepared runtime state 的算子不是某个样例的特殊逻辑。契约中对应字段为runtime_wrapper含source、entry、tensor_list_descriptor_strategy: prepared_workspace_tail、prepare_entry、descriptor_bytes、descriptor_order解析实现见 operator_codegen.py。八、自动修复与模板扩展点自动修复项apply-remediation对静态检查 findings 应用机器可修复项支持四种 kind定义见 sk_codegen_lib.pyrename-symbol符号重命名如历史 sysArgs API 名 → 当前 API 名。remove-line-containing删除包含指定内容的行。add-include补充缺失头文件。replace-pattern模式替换。不可自动修复项会作为人工处理项输出。新自动修复规则的扩展方式是向AUTO_REMEDIATION_KINDS增加 kind并在apply_remediation中增加处理分支。模板扩展点新基础算子模板新增templates/id.yaml。仓库自带的 add_custom.yaml 是一个最小非 SK elementwise add 算子模板渲染出一个干净的__global__ __vector__kernel含dtype参数可选float16/float32/int32可直接作为adapt-sk-from-global的干净输入演示闭环。新自动修复规则扩展scripts/sk_codegen_lib.py中的AUTO_REMEDIATION_KINDS。何时直接运行本工具只想确认一个算子是否能被识别detect-sk-form。生成阶段失败需要单独重跑adapt-sk-from-global。想检查聚合后的目录是否满足后续 pybind/wheel 构建要求。需要用apply-remediation对静态检查结果做自动修复尝试。验证命令python3 skills_root/sk-operator-codegen/scripts/operator_codegen.py --help九、在 skill 流水线中的位置与交付边界本 skill 的输出交给下游三个 skill 消费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 run-sk-pipeline编排完整闭环。端到端场景优先使用sk-operator-pipeline run-sk-pipelinesk-operator-codegen适合单独定位生成阶段问题。完整命令与产物说明见 SKILL.md 与 README.md。此外intake、plan、analyze-sk-conversion、adapt-sk-binding-scaffold、generate-sk-source-scaffold等历史 scaffold 入口仍保留供兼容场景使用。十、总结本手册作为本地行为契约sk-adaptation-cookbook.md之所以被定义为本地行为契约是因为它把 adapter 的行为约束到可测试、可复现的程度能自动生成的形态严格生成、不能确定的形态明确上报同时以原始__global__函数保持不变不猜变量名不静默回退 arch等硬规则保护生成产物的正确性。完整规格和边界情况以脚本实现、测试和本手册三者共同约束其中 sk_codegen_lib.py 是适配渲染器Args struct __sk__template SK_BIND与自动修复器的实现主体operator_codegen.py 是 CLI 入口与 IO 契约解析器二者共同构成本文所有规则的可执行依据。【免费下载链接】graph-autofusionGraph-autofusion 是一个面向昇腾Ascend芯片的轻量级、解耦式组件集合旨在通过自动融合技术加速模型执行。 目前已开源 SuperKernel 组件和 Autofuse 组件未来将持续开放更多自动融合相关模块。项目地址: https://gitcode.com/cann/graph-autofusion创作声明：本文部分内容由AI辅助生成（AIGC），仅供参考