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 自动适配为 SuperKernel（SK）binding 形态：从 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 struct：kernel 参数的结构化封装

命名为<NameCamel>Args，每个 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__函数

template<uint32_t splitidx> __sk__ <kernel_type> void <name>_sk(const <NameCamel>Args *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>_sk<0>, <name>_sk<1>, <name>_sk<2>, <name>_sk<3>)

SK_BIND将原始算子在 SK runtime 侧的 bind target 与若干 split 符号绑定：

• mask默认是4（DCCI，即 Dump/Compare/Check/Inspect 类能力 bit）。允许值是0..7，其中0表示没有能力 bit，1/2/4是 bit flag（可组合）。代码层面的取值范围校验在 sk_codegen_lib.py（mask must be in 0..7）。
• --num-splits控制绑定多少个<name>_sk<N>符号，范围1..4，校验在 sk_codegen_lib.py。CLI 默认值为4（见 operator_codegen.py）。
• 生成的宏文本形如SK_BIND(<bind_symbol>, <mask>, <split_args>);，其中<split_args>为name_sk<0>, name_sk<1>, ...的逗号拼接，见 sk_codegen_lib.py。

二、多算子聚合渲染：从单算子树到聚合 wheel

单算子适配产物（aclgraph-canonical 布局）

单算子适配仍然为每个 asset 写一个 aclgraph-canonical tree：

operator-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：

聚合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_ARCHS=dav-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：

形态判定的源码依据在 sk_codegen_lib.py：current-sk-bind要求同时存在__sk__与SK_BIND模板风格；partial指混合信号（例如有__sk__但没有SK_BIND）。operator_codegen.py中的_collect_sk_markers（operator_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）：

映射规则要点：

• __mix__特殊情形优先：当__mix__(c, v)中c >= 1 && v == 0时归一化为__cube__，当c == 0 && v >= 1时归一化为__vector__；其余__mix__(c, v)保持原样。
• 无法从 qualifier 推导出任何已知类型时，函数抛出ValueError（cannot map kernel type from qualifiers），避免生成非法 SK 签名。

五、sysArgs注入时机与 API 命名约束

三种模式

实现上sys_args_mode的三种取值判定在 sk_codegen_lib.py：auto模式直接读取解析阶段得到的entry.uses_get_block_num标志。

注入后的 API 命名

注入后必须使用当前 API 名：

• sysArgs->skNumBlocks
• sysArgs->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 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

生成 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 name=value \ --output-dir build/examples/sk-codegen/template-out

standalone 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-contract，adapt-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.py）；parameters还支持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.py）：

• rename-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-pipeline；sk-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