CANN graph-autofusion Codegen 代码生成模块深度解析从 ImplGraph 到可执行 Kernel 的完整链路【免费下载链接】graph-autofusionGraph-autofusion 是一个面向昇腾Ascend芯片的轻量级、解耦式组件集合旨在通过自动融合技术加速模型执行。 目前已开源 SuperKernel 组件和 Autofuse 组件未来将持续开放更多自动融合相关模块。项目地址: https://gitcode.com/cann/graph-autofusionCodegen 是 CANN graph-autofusion 编译流水线中的代码生成环节负责将调度阶段产出的 ImplGraph 解析为 Host 侧 tiling 函数与 Device 侧核函数源码实现一处调度、多处适配的多模板 kernel 分发。本文基于仓库中 Codegen 设计文档并结合 autofuse/codegen 与 autofuse/v35/codegen 下的源码实现完整梳理 Codegen 的设计思想、模块间契约、对外接口、分层架构与核心实现细节帮助读者理解融合算子自动代码生成的内部机制并掌握 ApiCall 扩展点与 v35 平台扩展层的接入方式。1. 特性背景与设计思想1.1 Codegen 在 Autofuse 流水线中的位置Autofuse 的编译流水线为图优化Optimize→ 调度Schedule→ 代码生成Codegen→ 性能调优ATT/Auto Tiling。各环节职责如下Schedule负责对融合后的计算图做切分决策UB 切分、核间切分、多模板输出 ImplGraphCodegen解析 ImplGraph生成实际可执行的代码包括在 Host 侧Host Code和 Device 侧Kernel Code运行的代码ATT负责在运行前计算 tiling 的具体取值tiling_key、各切分轴 size 等填充到 Codegen 生成的数据结构中。Codegen 的核心设计思想是Codegen 只生成代码的形不决定 tiling 的值。代码中所有与切分相关的量tiling_key、轴 size、循环次数都以 tiling_data 变量的形式占位具体数值由 ATT 在 Host 侧运行时计算并传入从而使同一份生成代码可以适配不同的 shape 和硬件资源条件。这一设计在源码中有直接体现Codegen::Generate的入参是::ascir::FusedScheduledResult产出CodegenResultproto / tiling_data / tiling / kernel / infer_shape 五类字符串见 autofuse/codegen/codegen.h而CodegenOptions中的tiling_lib_path与tiling_lib_codegen_symbol进一步表明 tiling 代码的生成能力来自外部库Codegen 只负责组装。1.2 设计目标一处调度多处适配Schedule 针对不同切分策略生成多个 ImplGraph 模板Codegen 将其全部生成并通过 tiling_key 分发运行时按实际 shape 选择最优模板生成代码自包含kernel 源文件内联所有依赖的 AscendC API 源码可独立编译平台可扩展通过 ApiCall 注册机制与平台扩展层支持不同计算类别Vector/Cube/MicroAPI和不同 SoC 版本的算子接入Kernel 生成主流程保持平台无关。2. 用户使用场景场景说明静态 shape 融合编译TF/ATC 前端接入shape 编译期已知Codegen 生成寄存器 tiling_key 分发的多模板 kernel运行时零解析开销动态 shape 融合编译shape 编译期未知tiling 在 Host 侧运行时计算Codegen 生成运行时解析 tiling_data 的分发结构PyTorch Inductor 接入通过GenerateForInductor路径生成 Inductor 专用接口支持静态 shape 全 const tiling 前移和 TopN tiling 候选选择CVCube/Vector融合Matmul 等算子与 Vector 算子融合Codegen 生成多 group CV kernel 与配套 tiling wrapperPGO tiling 搜索生成独立的 tiling 搜索可执行程序在真实硬件上实测各 tiling 候选性能保存最优结果供编译期使用从入口层看上述场景分别对应 codegen.h 中的Generate通用路径、GenerateForInductorInductor 路径、GenerateTilingForInductor、GeneratorPgoPGO 路径等接口每种场景都有独立的生成入口与参数形态。3. 特殊背景及限制与相邻组件的隐式契约Codegen 与相邻组件协作时存在大量约束这些约束大部分无法从单个模块的代码中直接看出是模块间协作的隐式契约修改任何一侧都需要同步核对。3.1 与 Schedule 的契约ImplGraphImplGraph 是 Schedule 与 Codegen 之间唯一的接口Codegen 不做任何调度决策隐含约束包括轴完备性ImplGraph 上的轴必须构成完整的切分树ORIGINAL → TILE/BLOCK 的 OUT/IN 配对。Codegen 按 BLOCK_OUT核间并行不生成 for、BLOCK_IN生成核内偏移、其余轴生成 for 循环三类处理缺失配对会导致生成的循环结构不完整节点信息完备性每个节点的 repeats/strides/vectorized_strides 必须与 axis 维度一致dtype 组合必须在 API 支持范围内Codegen 生成前会做校验CheckGraphValidity校验失败直接终止编译。当前 Broadcast 的 int64 dtype 暂不在校验范围内豁免项表达式可翻译性图上的 shape/offset 表达式ge::Expression必须能翻译为tiling 变量 循环变量的组合。Codegen 不支持运行时无法求值的表达式例如依赖 Host 侧随机性的表达式无法出现在轴定义中。CheckGraphValidity的实现位于 autofuse/codegen/codegen_graph_check.cpp由 autofuse/codegen/codegen_kernel.cpp 在解析图之前调用。从源码看校验逻辑包括通过CollectInputDtypes按 IR 输入类型必选/动态收集节点输入 dtype 组合并校验 API 支持范围、校验 axis/repeats/strides/vectorized_strides 的维度一致性、逐节点调用 IR codegen 实现的IsNodeValid。3.2 与 ATT/Auto Tiling 的职责边界Codegen 决定结构ATT 决定取值AutofuseTilingData的字段哪些轴需要 size 变量、tiling_key、block_dim由 Codegen 根据切分策略确定但字段的具体数值由 ATT 计算。两侧对 TilingData 结构的定义必须一致该一致性由两者共用同一份 TilingData 生成代码保证任何一侧单独修改字段都会导致运行时错位性能建模在 ATT 侧Codegen 在组装 API 参数时会计算outer_call_countAPI 外层调用次数等信息写入节点参数Auto Tiling 的性能公式耗时 ≈ API 调用次数 × 单次耗时以此为输入做 tiling 候选评估。性能公式本体不在 Codegen 中Codegen 新增 API 参数时需确认 ATT 侧公式是否依赖该参数PGO/TopN 的兜底限制tiling_key 数量超过上限10000时 PGO 搜索自动关闭回退到常规 tiling 生成tiling_key 数量超过单编译单元上限时kernel 会自动拆分为多个编译单元分别编译两处上限不同独立生效。3.3 与外部 tiling 库的依赖Host 侧 tiling 函数的实际生成逻辑位于外部 gen_tiling 库中Codegen 通过函数指针TilingLibCodegenFunc指向 gen_tiling.so 的CodegenTiling符号回调生成再由 Codegen 负责组装头文件、拼接翻译单元。这意味着tiling 代码生成能力的变更可能涉及 gen_tiling 库的同步更新两者需配套发布gen_tiling 库缺失或符号不匹配时tiling 生成失败与 kernel 生成是独立的两条失败路径。CodegenOptions::tiling_lib_path与tiling_lib_codegen_symbol两个配置项即对应上述回调机制前者指定 gen_tiling 库路径后者指定库内符号名见 autofuse/codegen/codegen.h。3.4 平台SoC 版本约束不同 SoC 版本如 3510、5102、9202的指令集和 API 形态存在差异版本差异通过 IR 注册表按 SoC 版本路由到不同的实现V2 实现Codegen 主流程不感知具体芯片新增 SoC 版本支持时需同时提供该版本的 ApiCall 实现和 API 源码模板注册不能复用其他版本的模板Kernel 文件中调用的 AscendC API 源码由AscendCApiRegistry在构建期从源码头文件打包为字符串注册生成时内联。API 模板更新后必须重新编译 Codegen 模块否则生成结果仍是旧模板。3.5 其他限制IO 数量限制单个融合算子的输入输出数超过 64 时kernel 形参自动切换为 list_tensor 形式超出该数量可能触发编译失败队列深度与 UB 容量UB 内存分配队列深度、double buffer由 ImplGraph 上的 mem/que 信息给出Codegen 忠实翻译不做事后调整UB 超限的校验发生在 Schedule 阶段Cube 算子的 tiling 依赖官方实现Matmul/Conv 的 tiling 通过 wrapper 调用官方 matmul tiling 接口AutofuseDoCubeMatMulTilingCodegen 不自行实现 Cube tiling 算法。list_tensor 机制在 autofuse/codegen/codegen_kernel.h 的Kernel::SetUseListTensor与KernelFuncDeclare(..., bool use_list_tensor)等接口中有对应实现Cube tiling 的 wrapper 调用则对应GenCubeTilingFuncCall、GenCubeCommonTiling等成员函数codegen_kernel.h。4. 对外接口4.1 生成产物接口运行时可见Codegen 产出的代码以 C 接口符号对外提供供运行时框架调用符号侧别说明AutofuseTilingGetTiling 系列Hosttiling 主函数根据 shape 计算 tiling_key、各轴切分大小、workspace 大小InferShapeHost形状推断函数输入符号表达式上下文输出各 tensor 形状GetKernelBin(std::vectorchar)Host返回已编译 kernel 的二进制内容供运行时直接加载免去运行期编译核函数extern C __global__ __aicore__Device按 tiling_key 分发多个模板函数的核函数入口AutofuseTilingData共享Host/Device 共享的 tiling 数据结构字段值由 ATT 填充4.2 内部扩展接口编译期接口说明Codegen::Generate / GenerateForInductor总入口输入 Schedule 的FusedScheduledResult输出CodegenResultproto/tiling_data/tiling/kernel/infer_shapeCodegen::GenerateTilingData / GenerateTiling / GenerateKernel / GenerateInferShape / GenGetKernelAndJson各产物的独立生成接口可单独调用ApiCallRegisterTApiCall 静态自注册新增算子只需实现ApiCall子类并注册主流程自动发现MicroApiCallRegisterTMicroAPI 子工厂注册v35 平台扩展层内部使用AscendCApiRegistry::RegisterApi / GetFileContentAPI 源码模板注册与查询生成 kernel 时按节点依赖内联CodegenResult结构体proto、tiling_data、tiling、kernel、infer_shape 五个字符串字段与全部独立生成接口的声明均位于 autofuse/codegen/codegen.h可作为二次开发时查阅接口清单的第一手资料。5. 整体架构Codegen 由门面类Codegenautofuse/codegen/codegen.h对外提供统一入口一次生成产出五类结果CodegenResult产物说明生成入口proto序列化的调度结果等元信息Codegen::Generatetiling_dataAutofuseTilingData结构体定义Codegen::GenerateTilingDatatilingHost 侧 tiling 函数源码Codegen::GenerateTilingkernelDevice 侧核函数源码Codegen::GenerateKernelinfer_shapeHost 侧形状推断函数源码Codegen::GenerateInferShape整体架构分层如下┌─────────────────────────────────────────────────────────────┐ │ 入口层Codegen 门面 │ │ Generate / GenerateForInductor / GenerateTiling / ... │ ├─────────────────────────────────────────────────────────────┤ │ Kernel 生成层 │ │ Kernel图解析 → Loop 树 → 循环与 API 调用生成 │ │ tiling_key 多模板分发 / API 源码内联 │ ├─────────────────────────────────────────────────────────────┤ │ ApiCall 框架核心扩展点 │ │ ApiCall 基类 ApiCallFactory 注册工厂 │ │ ┌──────────────┐ ┌─────────────────────────────────────┐ │ │ │ 主框架内置 │ │ v35 平台扩展层 │ │ │ │ Load/Store/ │ │ reg_api_callRegBase │ │ │ │ elewise/ │ │ cube_api_callMatmul/Conv2D │ │ │ │ reduce/ │ │ micro_api_callMicroAPI │ │ │ │ transpose... │ │ vec_func_callVectorFunc │ │ │ └──────────────┘ └─────────────────────────────────────┘ │ ├─────────────────────────────────────────────────────────────┤ │ Host 生成层 │ │ TilingData / TilingLib含 PGO、Inductor TopN、CV tiling │ │ InfershapeGen / GenGetKernelAndJson │ ├─────────────────────────────────────────────────────────────┤ │ 支撑层 │ │ AscendCApiRegistryAPI 模板内联 │ │ ExpressionConvertStruct表达式翻译 │ │ CheckGraphValidity图校验 / ApiParamDump调试 │ └─────────────────────────────────────────────────────────────┘生成顺序为先 Kernel再 TilingData、Tiling、InferShape最后可选执行 PGO 生成。Tiling 部分的实际生成由TilingLib通过函数指针回调外部 gen_tiling 库完成Codegen 负责组装与拼接各翻译单元。6. 核心实现6.1 输入ImplGraphCodegen 的输入是 Schedule 生成的 ImplGraph该图包含了 shape 信息、节点信息、轴信息等生成代码的基本要素。图上的具体信息如下// size描述计算量大小 Sizes: z0z1t_size: VAR z0z1Tb_size: VAR // axis描述轴信息包括大小、类型、对齐方式等 Axis: z0(0) : 200, ORIGINAL, align: -1, allow_oversize_axis: 0, allow_unaligned_tail: 1 z1(1) : 200, ORIGINAL, align: -1, allow_oversize_axis: 0, allow_unaligned_tail: 1 z0z1(2) : 40000, ORIGINAL, align: -1, allow_oversize_axis: 0, allow_unaligned_tail: 1 // ORIGINAL代表未切分的原始轴 z0z1T(3) : Ceiling((40000 / (z0z1t_size))), TILE_OUT, from: {z0z1, }, align: 1, allow_oversize_axis: 0, allow_unaligned_tail: 1 z0z1t(4) : z0z1t_size, TILE_IN, from: {z0z1, }, align: 1, allow_oversize_axis: 0, allow_unaligned_tail: 1 // TILE_IN代表UB切分的内轴 z0z1TB(5) : Ceiling((Ceiling((40000 / (z0z1t_size))) / (z0z1Tb_size))), BLOCK_OUT, from: {z0z1T, }, align: 1, allow_oversize_axis: 0, allow_unaligned_tail: 1 z0z1Tb(6) : z0z1Tb_size, BLOCK_IN, from: {z0z1T, }, align: 1, allow_oversize_axis: 0, allow_unaligned_tail: 1 // BLOCK_IN代表核间切分的内轴 // node每个节点会生成kernel代码中的api调用 Nodes: ...... abs_test/gather_0: Load (1) ...... abs_test/abs_0: Abs (2) .axis {z0z1TB, z0z1Tb, z0z1t, } .loop_axis z0z1Tb .api: .compute_type elewise .type Compute .unit Vector .x abs_test/gather_0.y .y.dtype float32 .y.axis {z0z1TB, z0z1Tb, z0z1t, } .y.repeats {(40000 / (z0z1Tb_size * z0z1t_size)), z0z1Tb_size, z0z1t_size, } .y.strides {(z0z1Tb_size * z0z1t_size), z0z1t_size, 1, } .y.vectorized_axis {z0z1t, } // 向量化轴代表UB内计算的数据量对应的轴 .y.vectorized_strides {1, } .y.mem: .tensor_id 3 .alloc_type Queue .hardware UB .position TPosition::VECOUT .y.que: .id 1 .depth 2 .buf_num 2 .reuse_id 1 abs_test/store: Store (3) ......从上述结构可以看出ImplGraph 以三种信息支撑代码生成Sizes计算量大小通常为 VAR 占位、Axis轴的类型 ORIGINAL/TILE_OUT/TILE_IN/BLOCK_OUT/BLOCK_IN、大小表达式、对齐与尾部处理属性、Nodes每个节点含 axis、repeats、strides、vectorized 信息与 mem/que 的 UB 内存分配信息。Axis、Tensor、Tiler、TPipe等解析期数据模型类均定义在 autofuse/codegen/codegen_kernel.hTiler负责将轴注册为 tiling 数据变量映射并提供Size/ActualSize/Offset等表达式求值接口TPipe负责 tensor/队列/缓冲区TQue/TBuf/MergeScope的登记与 UB 内存分配代码生成。根据 ImplGraph 解析以下示例代码展示了 Codegen 生成 DataCopy API 的过程可以看出核心流程包括解析图上的切分策略组装 API 的入参生成在 Device 侧执行的代码ss DataCopyPadExtend( ub , gm [ gm_offset tpipe.tiler.Size(api_attr.offset) ], dma_param.block_count , dma_param.block_len , dma_param.src_stride , dma_param.dst_stride ); std::endl;这里tpipe.tiler.Size(...)把轴偏移表达式翻译成 tiling 变量参与运算的 C 表达式dma_param则是对齐后的搬运参数block_count/block_len/src_stride/dst_stride最终拼接出完整的数据搬运调用。6.2 Kernel 代码生成Kernel 生成由Kernel类autofuse/codegen/codegen_kernel.h完成一张 ImplGraph 对应一个 Kernel 生成器整体分为解析和生成两个阶段。6.2.1 图解析与合法性校验Kernel::ParseGraph解析 ImplGraph 前首先调用CheckGraphValidityautofuse/codegen/codegen_graph_check.cpp对图做合法性校验包括dtype 校验收集节点输入 dtype 组合校验 API 是否支持repeats/strides 校验校验 axis、repeats、strides、vectorized_strides 维度一致性节点校验逐节点调用 IR codegen 实现的IsNodeValid。解析过程中Kernel 将图上的轴注册进Tilertiling 数据变量映射、将 tensor/队列注册进TPipeUB 内存与队列管理并按节点序列构造出Loop树autofuse/codegen/codegen_kernel_loop.cpp树中每个节点对应一个 ApiCall 对象。Kernel::ParseGraph的静态签名定义于 codegen_kernel.h。6.2.2 循环与 API 调用生成Loop::Generate按 ImplGraph 的切分策略递归生成循环结构BLOCK_OUT 轴核间并行轴只生成 size/actual_size 计算不生成 for 循环通过 blockidx 隐式并行BLOCK_IN 轴生成block_dim_offset偏移计算确定当前核处理的数据分片其余轴生成for (axis 0; axis loop_size; axis)循环并在循环体内生成尾块tail的实际 size 计算与缓存使能条件。循环体内每遇到一个计算节点按固定骨架生成代码WaitInputs → AllocOutputs → API 调用 → SyncOutputs → FreeInputs → FreeUnusedOutputs即围绕 API 调用完成 UB 队列的等待、分配、同步与释放保证 MTE 搬运与 Vector 计算的正确流水。6.2.3 多模板 tiling_key 分发与传统算子代码类似针对不同的切分策略Schedule 会生成多个模板对应不同的 ImplGraphCodegen 需要依次解析这些模板生成模板函数并在核函数入口处根据 tiling_key 调用不同的模板函数。示例如下extern C __global__ __aicore__ void abs_test(GM_ADDR abs_test_Data_0, GM_ADDR abs_test_Output_0, GM_ADDR workspace, AutofuseTilingData param) { const AutofuseTilingData t; if (t.tiling_key 0) { abs_test_0_general_0_nil_2_nil(abs_test_Data_0, abs_test_Output_0, workspace, t); } else if (t.tiling_key 1) { abs_test_0_general_0_nil_2_nil_unaligned(abs_test_Data_0, abs_test_Output_0, workspace, t); } }上面AutofuseTilingData结构体示例如下其中的成员变量具体值包括 tiling_key 等由 ATT 计算给出BEGIN_TILING_DATA_DEF_T(AutofuseTilingData) const uint32_t block_dim 40; const uint32_t corenum 0; const uint32_t ub_size 196352; const uint32_t hbm_size 0; const uint32_t tiling_key 0; const uint32_t z0z1z2z3t_size 17968; const uint32_t z0z1z2z3Tb_size 57; const uint32_t q0_size 96; const uint32_t q1_size 35936; const uint32_t q2_size 35936; const uint32_t b0_size 35936; const uint32_t tmp_tbuf_size 8192; END_TILING_DATA_DEF_T;不同的模板函数生成的 for 循环以及每次处理的数据量各不相同示例如下tiling_key0对 z1 轴进行切分每次在 UB 中完成计算的向量化轴为 z1t。for (int z0z1Tb 0; z0z1Tb z0z1Tb_loop_size; z0z1Tb) { Abs(y_local[0], xlocal[0], z1t_actual_size); }tiling_key1对 z0 轴进行切分每次在 UB 中完成计算的向量化轴为 z0t 以及 z1。for (int z0Tb 0; z0Tb z0Tb_loop_size; z0Tb) { Abs(y_local[0], xlocal[0], z0t_actual_size * z1_axis_size); }Kernel::GenKernelFuncByTilingKey根据场景选择不同的分发实现对应 codegen_kernel.h 声明寄存器 tiling_key静态 shape将 tiling_key 写入寄存器每个调度模板生成一个 kernel 函数入口处生成if (t.tiling_key N) { ... }分发链当 group 内多模板需要并行时插入SyncAll()同步。其具体实现细分为GenKernelFuncWithRegTilingKey、GenSingleGroupKernelWithRegTilingKey、GenMulGroupKernelWithRegTilingKey三个静态方法解析 tiling_data动态 shape / Inductor运行时解析 tiling_data 结构选择分支对应GenKernelFuncWithParseTilingData及其单/多 group 变体CV 融合生成多 group 的 CV kernel 函数对应GenCVKernelFuncWithMulGroup模板数超限tiling_key 数量超过单编译单元上限时自动拆分为多个编译单元PackingFunc分别编译对应GenPackingFunctions/GetMaxGroupPerCompileUnit等私有方法。6.2.4 API 源码内联Kernel 文件中调用的 AscendC API如DataCopyPadExtend、Abs等的实现源码由AscendCApiRegistryautofuse/codegen/ascendc_api_registry.cpp统一管理。构建期将各 API 头文件打包为源码字符串注册进单例表生成 kernel 时按节点声明的头文件依赖将 API 源码字符串直接内联到 kernel 文件中使每个 kernel 源文件自包含、可独立编译。6.3 Host Code 生成6.3.1 tiling_data 生成TilingData::Generateautofuse/codegen/codegen_tiling_data.cpp生成AutofuseTilingData结构体定义字段包括 tiling_key、block_dim、workspace 大小、各切分轴 size 等多 ScheduleGroup 场景生成统一的包装结构。对于静态 shape 的 Inductor 场景还会生成全 const 初值版本将 tiling 计算前移到编译期。6.3.2 tiling_func 生成Host 侧 tiling 函数由TilingLib组织生成autofuse/codegen/codegen_tiling.cpp包括AutofuseGetTilingSize估算 tiling 所需内存、AutofuseTiling主 tiling 函数根据 shape 计算 tiling_key 和各轴切分大小、workspace 计算等函数。除常规生成外tiling 侧还支持PGO 生成生成独立的 tiling 搜索可执行程序在真实硬件上遍历 tiling 候选、采集性能并保存最优结果Inductor TopN 选择结合性能建模公式由 ATT 提供耗时 ≈ API 调用次数 × 单次耗时与实测 profiling对 tiling 候选做建模评估与 TopN 筛选CV 融合 tiling针对 Cube/Vector 融合场景生成专用 tiling 函数与辅助接口。PGO 相关实现进一步拆分在 autofuse/codegen/codegen_tiling_pgo_search.cpp、codegen_tiling_pgo_memory.cpp、codegen_tiling_pgo_runtime.cpp 等文件中GeneratorPgo入口声明见 codegen.h。6.3.3 infer_shape 生成InfershapeGen::GenInferShapeFuncautofuse/codegen/codegen_infershape.cpp生成extern C InferShape函数解析各输出 shape 表达式中的符号变量生成逐维 AppendDim 的推断代码对含除法取整的表达式额外生成四舍五入与精度校验逻辑保证动态 shape 下形状推断与 kernel 实际计算一致。6.3.4 get_kernel 生成Codegen::GenGetKernelAndJson读取已编译的 kernel 二进制ota生成extern C void GetKernelBin(std::vectorchar)函数将 Device 代码以二进制形式嵌入 Host 侧源码供运行时直接加载免去运行期编译开销。6.4 ApiCall 框架ImplGraph 上每个节点最终都会转化为一次 API 调用这一转换由 ApiCall 框架承载是 Codegen 的核心扩展点。6.4.1 基类与生成链路所有 ApiCall 继承自基类ApiCallautofuse/codegen/codegen_kernel_loop.h核心生成链路为Init(node) // 收集 unit、输入输出 tensor、tmp_buf 信息 → ParseAttr() // 子类覆写解析算子属性 → BuildApiParam() // 组装 CodegenApiParam循环轴、offset、cal_count 等 → GenerateApiCallString()// 按 api_param 拼接 API 调用字符串 → PostProcess() // 生成收尾同步代码其中CodegenApiParam是 API 入参的中间表示将图上的符号表达式ge::Expression经ExpressionConvertStructautofuse/codegen/expression_convert_struct.cpp翻译为 tiling 变量、循环变量的 C 表达式组合例如cal_count、偏移量、循环 size 等。基类还提供 UB 生命周期辅助接口WaitInputs/AllocOutputs/SyncOutputs/FreeInputs等处理共享队列复用与 MTE2/MTE3 同步子类通常只需覆写BuildApiParam和Generate即可接入新算子。6.4.2 工厂与注册机制ApiCall 的创建采用静态自注册工厂模式autofuse/codegen/api_call/utils/api_call_factory.h各 ApiCall 实现类在文件尾定义静态ApiCallRegisterT对象向单例工厂ApiCallFactory注册类名 → 构造函数IR 定义侧ascir codegen impl通过GetApiCallName()声明节点对应的 ApiCall 类名Kernel 生成遍历节点时按类名从工厂创建 ApiCall 对象CreateApiCallObject见 codegen_kernel_loop.cpp主生成流程无需感知具体算子实现新增算子只需增加实现类并注册。6.4.3 内置 ApiCall主框架autofuse/codegen/api_call/内置的 ApiCall 按计算类别组织类别ApiCall数据搬运LoadApiCall、StoreApiCall一元计算UnaryApiCall、CastApiCall、RsqrtApiCall、LogicalNotApiCall、UnaryBitWidthChangeApiCall等二元计算BinaryApiCall、AxpyApiCall、TrueDivApiCall、PowApiCall、LeakyReluApiCall等比较与选择CompareApiCall、WhereApiCall、ClipByValueApiCall归约ReduceApiCall数据重排TransposeApiCall、BroadcastApiCall、ConcatApiCall、GatherApiCall其他PadApiCall、RemovePadApiCall、Ub2UbApiCall等主框架内置实现的目录组织为 autofuse/codegen/api_call/ 下的 broadcast、concat、datacopy、elewise、gather、reduce、transpose、utils 等子目录与上表类别一一对应。6.5 v35 平台扩展层不同 SoC 版本的硬件指令集和 API 形态存在差异Codegen 通过平台扩展层autofuse/v35/codegen/承载版本特定实现。扩展层完全复用主框架的基类与注册机制各实现类同样继承ApiCall并静态注册进ApiCallFactory由 SoC 版本路由在 IR 注册表层面按芯片版本选择实现Kernel 生成流程本身保持平台无关。v35 扩展层包含四类实现6.5.1 reg_api_callRegBase 风格 API面向 RegTensor 指令形态的向量计算与数据搬运 API是 v35 的主体包括BinaryApiCallV2、CastV2ApiCall、CompareV2ApiCall、LoadRegApiCall、StoreRegApiCall、NddmaApiCall、IndirectLoadRegApiCall、BroadcastRegApiCall、ConcatRegApiCall、SplitRegApiCall、RegReduceApiCall、TransposeRegApiCall、GatherRegApiCall、WhereRegApiCall、SoftmaxApiCall等约 24 个类。相比主框架版本V2 系列 API 具备更强的参数化能力如轴合并 merge_axes并会在BuildApiParam中计算outer_call_countAPI 外层调用次数该值直接对接 Auto Tiling 的性能公式建模耗时 ≈ 调用次数 × 单次耗时作为 tiling 候选评估的输入。对应的 API 源码模板约 100 个 RegBase API由构建期从源码头文件自动打包为字符串注册进AscendCApiRegistryautofuse/v35/codegen/ascendc_reg_base_api_register.cpp。6.5.2 cube_api_callCube 算子 APIMatmulApiCall、Conv2DApiCall生成 Matmul/Conv 等 Cube 算子的 AscendC API 调用。Cube 单元调用除生成调用点外还会生成模板函数定义GenerateFuncDefinition和宏展开GenerateMacro并通过 tiling wrapper 对接官方 matmul tiling 实现。相关 Cube API 模板matmul、batch_matmul、conv2d 及其 dynamic/tiling_key/pingpong 变体注册见 autofuse/v35/codegen/ascendc_cube_api_register.cpp。6.5.3 micro_api_callMicroAPIMicroAPI 是最低一级的代码形态直接生成AscendC::MicroAPI::*指令级调用如LoadAlign、StoreAlign与AscendC::Reg::*寄存器操作如Pack、UnPack、MaskPack由独立的MicroApiCall基类与MicroApiCallFactory工厂承载不继承主框架ApiCall包括MicroLoadApiCall、MicroStoreApiCall、MicroCastApiCall、MicroCompareApiCall、MicroWhereApiCall、MicroBinaryScalarApiCall、MicroScalarBroadcastApiCall等。6.5.4 vec_func_callVectorFuncVfCall将 ImplGraph 中的子图封装为一个独立的 vf 函数实现多次循环 → 一次函数调用的结构优化。其内部持有VFLoop循环树遍历子图节点时通过 MicroAPI 工厂为每个节点创建 MicroApiCall形成两层结构VfCall主框架 ApiCall生成函数定义与调用点 └── VFLoopvf 内循环树 └── MicroApiCall指令级 API 调用该机制同时支持 CV-UB 融合场景的 kernel 生成与Kernel::GenerateVecFuncOfCVFusion、InitCVFusionAddr等 CV 融合相关实现codegen_kernel.h配套工作。6.6 辅助机制API 参数 dumpCodegenApiParam::DumpGraphApiParamsautofuse/codegen/codegen_api_param_dump.cpp可将每张 ImplGraph 中所有节点的 API 入参序列化为文本文件用于调试和对比生成的调用参数开发新 ApiCall 时可通过该工具核验参数组装是否正确表达式转换ExpressionItem/CombinedExpression提供符号表达式到 C 代码的统一渲染并附带工厂函数ExprItemFactory/CombinedExprFactory用于组装常见表达式模式如SizeMinusActualSize尾块计算。7. 关键文件索引文件路径职责autofuse/codegen/codegen.h / codegen.cppCodegen 门面类总入口与产物组装autofuse/codegen/codegen_kernel.h / codegen_kernel.cppKernel 生成器、Tiler/TPipe 数据模型、tiling_key 分发autofuse/codegen/codegen_kernel_loop.h / codegen_kernel_loop.cppLoop 循环树、ApiCall 基类与生成骨架autofuse/codegen/codegen_tiling.cpp 等Host 侧 tiling 生成含 PGO、Inductor TopN、CV tilingautofuse/codegen/codegen_tiling_data.cppAutofuseTilingData 结构体生成autofuse/codegen/codegen_infershape.cppInferShape 函数生成autofuse/codegen/codegen_graph_check.cppImplGraph 合法性校验autofuse/codegen/expression_convert_struct.cpp符号表达式到 C 表达式转换autofuse/codegen/codegen_api_param_dump.cppAPI 入参 dump 调试工具autofuse/codegen/ascendc_api_registry.cppAscendC API 源码模板注册表autofuse/codegen/api_call/主框架内置 ApiCall 实现datacopy/elewise/reduce/transpose/gather/broadcast/concatautofuse/v35/codegen/reg_api_call/RegBase 风格 V2 ApiCall平台扩展层主体autofuse/v35/codegen/cube_api_call/Matmul/Conv2D Cube ApiCallautofuse/v35/codegen/micro_api_call/MicroAPI 指令级 ApiCall 及独立工厂autofuse/v35/codegen/vec_func_call/VectorFunc 子图封装VfCall/VFLoopautofuse/v35/codegen/ascendc_reg_base_api_register.cppRegBase API 模板注册autofuse/v35/codegen/ascendc_cube_api_register.cppCube API 模板注册8. 小结Codegen 以生成代码的形、不决定 tiling 的值为根本设计原则通过 ImplGraph 这一唯一接口与 Schedule 解耦通过AutofuseTilingData与 ATT 分工通过 ApiCall 静态注册工厂实现算子接入的平台无关扩展再以 v35 平台扩展层承载不同 SoC 版本的指令形态差异。理解这一套契约 门面 扩展点的架构是向 graph-autofusion 中新增融合算子、适配新 SoC 版本或扩展 tiling 策略的起点而 Codegen 设计文档 本身则是最佳的架构导览本文所列源码路径可作为逐模块深入阅读的索引。【免费下载链接】graph-autofusionGraph-autofusion 是一个面向昇腾Ascend芯片的轻量级、解耦式组件集合旨在通过自动融合技术加速模型执行。 目前已开源 SuperKernel 组件和 Autofuse 组件未来将持续开放更多自动融合相关模块。项目地址: https://gitcode.com/cann/graph-autofusion创作声明：本文部分内容由AI辅助生成（AIGC），仅供参考 CANN graph-autofusion Codegen 代码生成模块深度解析：从 ImplGraph 到可执行 Kernel 的完整链路
CANN graph-autofusion Codegen 代码生成模块深度解析从 ImplGraph 到可执行 Kernel 的完整链路【免费下载链接】graph-autofusionGraph-autofusion 是一个面向昇腾Ascend芯片的轻量级、解耦式组件集合旨在通过自动融合技术加速模型执行。 目前已开源 SuperKernel 组件和 Autofuse 组件未来将持续开放更多自动融合相关模块。项目地址: https://gitcode.com/cann/graph-autofusionCodegen 是 CANN graph-autofusion 编译流水线中的代码生成环节负责将调度阶段产出的 ImplGraph 解析为 Host 侧 tiling 函数与 Device 侧核函数源码实现一处调度、多处适配的多模板 kernel 分发。本文基于仓库中 Codegen 设计文档并结合 autofuse/codegen 与 autofuse/v35/codegen 下的源码实现完整梳理 Codegen 的设计思想、模块间契约、对外接口、分层架构与核心实现细节帮助读者理解融合算子自动代码生成的内部机制并掌握 ApiCall 扩展点与 v35 平台扩展层的接入方式。1. 特性背景与设计思想1.1 Codegen 在 Autofuse 流水线中的位置Autofuse 的编译流水线为图优化Optimize→ 调度Schedule→ 代码生成Codegen→ 性能调优ATT/Auto Tiling。各环节职责如下Schedule负责对融合后的计算图做切分决策UB 切分、核间切分、多模板输出 ImplGraphCodegen解析 ImplGraph生成实际可执行的代码包括在 Host 侧Host Code和 Device 侧Kernel Code运行的代码ATT负责在运行前计算 tiling 的具体取值tiling_key、各切分轴 size 等填充到 Codegen 生成的数据结构中。Codegen 的核心设计思想是Codegen 只生成代码的形不决定 tiling 的值。代码中所有与切分相关的量tiling_key、轴 size、循环次数都以 tiling_data 变量的形式占位具体数值由 ATT 在 Host 侧运行时计算并传入从而使同一份生成代码可以适配不同的 shape 和硬件资源条件。这一设计在源码中有直接体现Codegen::Generate的入参是::ascir::FusedScheduledResult产出CodegenResultproto / tiling_data / tiling / kernel / infer_shape 五类字符串见 autofuse/codegen/codegen.h而CodegenOptions中的tiling_lib_path与tiling_lib_codegen_symbol进一步表明 tiling 代码的生成能力来自外部库Codegen 只负责组装。1.2 设计目标一处调度多处适配Schedule 针对不同切分策略生成多个 ImplGraph 模板Codegen 将其全部生成并通过 tiling_key 分发运行时按实际 shape 选择最优模板生成代码自包含kernel 源文件内联所有依赖的 AscendC API 源码可独立编译平台可扩展通过 ApiCall 注册机制与平台扩展层支持不同计算类别Vector/Cube/MicroAPI和不同 SoC 版本的算子接入Kernel 生成主流程保持平台无关。2. 用户使用场景场景说明静态 shape 融合编译TF/ATC 前端接入shape 编译期已知Codegen 生成寄存器 tiling_key 分发的多模板 kernel运行时零解析开销动态 shape 融合编译shape 编译期未知tiling 在 Host 侧运行时计算Codegen 生成运行时解析 tiling_data 的分发结构PyTorch Inductor 接入通过GenerateForInductor路径生成 Inductor 专用接口支持静态 shape 全 const tiling 前移和 TopN tiling 候选选择CVCube/Vector融合Matmul 等算子与 Vector 算子融合Codegen 生成多 group CV kernel 与配套 tiling wrapperPGO tiling 搜索生成独立的 tiling 搜索可执行程序在真实硬件上实测各 tiling 候选性能保存最优结果供编译期使用从入口层看上述场景分别对应 codegen.h 中的Generate通用路径、GenerateForInductorInductor 路径、GenerateTilingForInductor、GeneratorPgoPGO 路径等接口每种场景都有独立的生成入口与参数形态。3. 特殊背景及限制与相邻组件的隐式契约Codegen 与相邻组件协作时存在大量约束这些约束大部分无法从单个模块的代码中直接看出是模块间协作的隐式契约修改任何一侧都需要同步核对。3.1 与 Schedule 的契约ImplGraphImplGraph 是 Schedule 与 Codegen 之间唯一的接口Codegen 不做任何调度决策隐含约束包括轴完备性ImplGraph 上的轴必须构成完整的切分树ORIGINAL → TILE/BLOCK 的 OUT/IN 配对。Codegen 按 BLOCK_OUT核间并行不生成 for、BLOCK_IN生成核内偏移、其余轴生成 for 循环三类处理缺失配对会导致生成的循环结构不完整节点信息完备性每个节点的 repeats/strides/vectorized_strides 必须与 axis 维度一致dtype 组合必须在 API 支持范围内Codegen 生成前会做校验CheckGraphValidity校验失败直接终止编译。当前 Broadcast 的 int64 dtype 暂不在校验范围内豁免项表达式可翻译性图上的 shape/offset 表达式ge::Expression必须能翻译为tiling 变量 循环变量的组合。Codegen 不支持运行时无法求值的表达式例如依赖 Host 侧随机性的表达式无法出现在轴定义中。CheckGraphValidity的实现位于 autofuse/codegen/codegen_graph_check.cpp由 autofuse/codegen/codegen_kernel.cpp 在解析图之前调用。从源码看校验逻辑包括通过CollectInputDtypes按 IR 输入类型必选/动态收集节点输入 dtype 组合并校验 API 支持范围、校验 axis/repeats/strides/vectorized_strides 的维度一致性、逐节点调用 IR codegen 实现的IsNodeValid。3.2 与 ATT/Auto Tiling 的职责边界Codegen 决定结构ATT 决定取值AutofuseTilingData的字段哪些轴需要 size 变量、tiling_key、block_dim由 Codegen 根据切分策略确定但字段的具体数值由 ATT 计算。两侧对 TilingData 结构的定义必须一致该一致性由两者共用同一份 TilingData 生成代码保证任何一侧单独修改字段都会导致运行时错位性能建模在 ATT 侧Codegen 在组装 API 参数时会计算outer_call_countAPI 外层调用次数等信息写入节点参数Auto Tiling 的性能公式耗时 ≈ API 调用次数 × 单次耗时以此为输入做 tiling 候选评估。性能公式本体不在 Codegen 中Codegen 新增 API 参数时需确认 ATT 侧公式是否依赖该参数PGO/TopN 的兜底限制tiling_key 数量超过上限10000时 PGO 搜索自动关闭回退到常规 tiling 生成tiling_key 数量超过单编译单元上限时kernel 会自动拆分为多个编译单元分别编译两处上限不同独立生效。3.3 与外部 tiling 库的依赖Host 侧 tiling 函数的实际生成逻辑位于外部 gen_tiling 库中Codegen 通过函数指针TilingLibCodegenFunc指向 gen_tiling.so 的CodegenTiling符号回调生成再由 Codegen 负责组装头文件、拼接翻译单元。这意味着tiling 代码生成能力的变更可能涉及 gen_tiling 库的同步更新两者需配套发布gen_tiling 库缺失或符号不匹配时tiling 生成失败与 kernel 生成是独立的两条失败路径。CodegenOptions::tiling_lib_path与tiling_lib_codegen_symbol两个配置项即对应上述回调机制前者指定 gen_tiling 库路径后者指定库内符号名见 autofuse/codegen/codegen.h。3.4 平台SoC 版本约束不同 SoC 版本如 3510、5102、9202的指令集和 API 形态存在差异版本差异通过 IR 注册表按 SoC 版本路由到不同的实现V2 实现Codegen 主流程不感知具体芯片新增 SoC 版本支持时需同时提供该版本的 ApiCall 实现和 API 源码模板注册不能复用其他版本的模板Kernel 文件中调用的 AscendC API 源码由AscendCApiRegistry在构建期从源码头文件打包为字符串注册生成时内联。API 模板更新后必须重新编译 Codegen 模块否则生成结果仍是旧模板。3.5 其他限制IO 数量限制单个融合算子的输入输出数超过 64 时kernel 形参自动切换为 list_tensor 形式超出该数量可能触发编译失败队列深度与 UB 容量UB 内存分配队列深度、double buffer由 ImplGraph 上的 mem/que 信息给出Codegen 忠实翻译不做事后调整UB 超限的校验发生在 Schedule 阶段Cube 算子的 tiling 依赖官方实现Matmul/Conv 的 tiling 通过 wrapper 调用官方 matmul tiling 接口AutofuseDoCubeMatMulTilingCodegen 不自行实现 Cube tiling 算法。list_tensor 机制在 autofuse/codegen/codegen_kernel.h 的Kernel::SetUseListTensor与KernelFuncDeclare(..., bool use_list_tensor)等接口中有对应实现Cube tiling 的 wrapper 调用则对应GenCubeTilingFuncCall、GenCubeCommonTiling等成员函数codegen_kernel.h。4. 对外接口4.1 生成产物接口运行时可见Codegen 产出的代码以 C 接口符号对外提供供运行时框架调用符号侧别说明AutofuseTilingGetTiling 系列Hosttiling 主函数根据 shape 计算 tiling_key、各轴切分大小、workspace 大小InferShapeHost形状推断函数输入符号表达式上下文输出各 tensor 形状GetKernelBin(std::vectorchar)Host返回已编译 kernel 的二进制内容供运行时直接加载免去运行期编译核函数extern C __global__ __aicore__Device按 tiling_key 分发多个模板函数的核函数入口AutofuseTilingData共享Host/Device 共享的 tiling 数据结构字段值由 ATT 填充4.2 内部扩展接口编译期接口说明Codegen::Generate / GenerateForInductor总入口输入 Schedule 的FusedScheduledResult输出CodegenResultproto/tiling_data/tiling/kernel/infer_shapeCodegen::GenerateTilingData / GenerateTiling / GenerateKernel / GenerateInferShape / GenGetKernelAndJson各产物的独立生成接口可单独调用ApiCallRegisterTApiCall 静态自注册新增算子只需实现ApiCall子类并注册主流程自动发现MicroApiCallRegisterTMicroAPI 子工厂注册v35 平台扩展层内部使用AscendCApiRegistry::RegisterApi / GetFileContentAPI 源码模板注册与查询生成 kernel 时按节点依赖内联CodegenResult结构体proto、tiling_data、tiling、kernel、infer_shape 五个字符串字段与全部独立生成接口的声明均位于 autofuse/codegen/codegen.h可作为二次开发时查阅接口清单的第一手资料。5. 整体架构Codegen 由门面类Codegenautofuse/codegen/codegen.h对外提供统一入口一次生成产出五类结果CodegenResult产物说明生成入口proto序列化的调度结果等元信息Codegen::Generatetiling_dataAutofuseTilingData结构体定义Codegen::GenerateTilingDatatilingHost 侧 tiling 函数源码Codegen::GenerateTilingkernelDevice 侧核函数源码Codegen::GenerateKernelinfer_shapeHost 侧形状推断函数源码Codegen::GenerateInferShape整体架构分层如下┌─────────────────────────────────────────────────────────────┐ │ 入口层Codegen 门面 │ │ Generate / GenerateForInductor / GenerateTiling / ... │ ├─────────────────────────────────────────────────────────────┤ │ Kernel 生成层 │ │ Kernel图解析 → Loop 树 → 循环与 API 调用生成 │ │ tiling_key 多模板分发 / API 源码内联 │ ├─────────────────────────────────────────────────────────────┤ │ ApiCall 框架核心扩展点 │ │ ApiCall 基类 ApiCallFactory 注册工厂 │ │ ┌──────────────┐ ┌─────────────────────────────────────┐ │ │ │ 主框架内置 │ │ v35 平台扩展层 │ │ │ │ Load/Store/ │ │ reg_api_callRegBase │ │ │ │ elewise/ │ │ cube_api_callMatmul/Conv2D │ │ │ │ reduce/ │ │ micro_api_callMicroAPI │ │ │ │ transpose... │ │ vec_func_callVectorFunc │ │ │ └──────────────┘ └─────────────────────────────────────┘ │ ├─────────────────────────────────────────────────────────────┤ │ Host 生成层 │ │ TilingData / TilingLib含 PGO、Inductor TopN、CV tiling │ │ InfershapeGen / GenGetKernelAndJson │ ├─────────────────────────────────────────────────────────────┤ │ 支撑层 │ │ AscendCApiRegistryAPI 模板内联 │ │ ExpressionConvertStruct表达式翻译 │ │ CheckGraphValidity图校验 / ApiParamDump调试 │ └─────────────────────────────────────────────────────────────┘生成顺序为先 Kernel再 TilingData、Tiling、InferShape最后可选执行 PGO 生成。Tiling 部分的实际生成由TilingLib通过函数指针回调外部 gen_tiling 库完成Codegen 负责组装与拼接各翻译单元。6. 核心实现6.1 输入ImplGraphCodegen 的输入是 Schedule 生成的 ImplGraph该图包含了 shape 信息、节点信息、轴信息等生成代码的基本要素。图上的具体信息如下// size描述计算量大小 Sizes: z0z1t_size: VAR z0z1Tb_size: VAR // axis描述轴信息包括大小、类型、对齐方式等 Axis: z0(0) : 200, ORIGINAL, align: -1, allow_oversize_axis: 0, allow_unaligned_tail: 1 z1(1) : 200, ORIGINAL, align: -1, allow_oversize_axis: 0, allow_unaligned_tail: 1 z0z1(2) : 40000, ORIGINAL, align: -1, allow_oversize_axis: 0, allow_unaligned_tail: 1 // ORIGINAL代表未切分的原始轴 z0z1T(3) : Ceiling((40000 / (z0z1t_size))), TILE_OUT, from: {z0z1, }, align: 1, allow_oversize_axis: 0, allow_unaligned_tail: 1 z0z1t(4) : z0z1t_size, TILE_IN, from: {z0z1, }, align: 1, allow_oversize_axis: 0, allow_unaligned_tail: 1 // TILE_IN代表UB切分的内轴 z0z1TB(5) : Ceiling((Ceiling((40000 / (z0z1t_size))) / (z0z1Tb_size))), BLOCK_OUT, from: {z0z1T, }, align: 1, allow_oversize_axis: 0, allow_unaligned_tail: 1 z0z1Tb(6) : z0z1Tb_size, BLOCK_IN, from: {z0z1T, }, align: 1, allow_oversize_axis: 0, allow_unaligned_tail: 1 // BLOCK_IN代表核间切分的内轴 // node每个节点会生成kernel代码中的api调用 Nodes: ...... abs_test/gather_0: Load (1) ...... abs_test/abs_0: Abs (2) .axis {z0z1TB, z0z1Tb, z0z1t, } .loop_axis z0z1Tb .api: .compute_type elewise .type Compute .unit Vector .x abs_test/gather_0.y .y.dtype float32 .y.axis {z0z1TB, z0z1Tb, z0z1t, } .y.repeats {(40000 / (z0z1Tb_size * z0z1t_size)), z0z1Tb_size, z0z1t_size, } .y.strides {(z0z1Tb_size * z0z1t_size), z0z1t_size, 1, } .y.vectorized_axis {z0z1t, } // 向量化轴代表UB内计算的数据量对应的轴 .y.vectorized_strides {1, } .y.mem: .tensor_id 3 .alloc_type Queue .hardware UB .position TPosition::VECOUT .y.que: .id 1 .depth 2 .buf_num 2 .reuse_id 1 abs_test/store: Store (3) ......从上述结构可以看出ImplGraph 以三种信息支撑代码生成Sizes计算量大小通常为 VAR 占位、Axis轴的类型 ORIGINAL/TILE_OUT/TILE_IN/BLOCK_OUT/BLOCK_IN、大小表达式、对齐与尾部处理属性、Nodes每个节点含 axis、repeats、strides、vectorized 信息与 mem/que 的 UB 内存分配信息。Axis、Tensor、Tiler、TPipe等解析期数据模型类均定义在 autofuse/codegen/codegen_kernel.hTiler负责将轴注册为 tiling 数据变量映射并提供Size/ActualSize/Offset等表达式求值接口TPipe负责 tensor/队列/缓冲区TQue/TBuf/MergeScope的登记与 UB 内存分配代码生成。根据 ImplGraph 解析以下示例代码展示了 Codegen 生成 DataCopy API 的过程可以看出核心流程包括解析图上的切分策略组装 API 的入参生成在 Device 侧执行的代码ss DataCopyPadExtend( ub , gm [ gm_offset tpipe.tiler.Size(api_attr.offset) ], dma_param.block_count , dma_param.block_len , dma_param.src_stride , dma_param.dst_stride ); std::endl;这里tpipe.tiler.Size(...)把轴偏移表达式翻译成 tiling 变量参与运算的 C 表达式dma_param则是对齐后的搬运参数block_count/block_len/src_stride/dst_stride最终拼接出完整的数据搬运调用。6.2 Kernel 代码生成Kernel 生成由Kernel类autofuse/codegen/codegen_kernel.h完成一张 ImplGraph 对应一个 Kernel 生成器整体分为解析和生成两个阶段。6.2.1 图解析与合法性校验Kernel::ParseGraph解析 ImplGraph 前首先调用CheckGraphValidityautofuse/codegen/codegen_graph_check.cpp对图做合法性校验包括dtype 校验收集节点输入 dtype 组合校验 API 是否支持repeats/strides 校验校验 axis、repeats、strides、vectorized_strides 维度一致性节点校验逐节点调用 IR codegen 实现的IsNodeValid。解析过程中Kernel 将图上的轴注册进Tilertiling 数据变量映射、将 tensor/队列注册进TPipeUB 内存与队列管理并按节点序列构造出Loop树autofuse/codegen/codegen_kernel_loop.cpp树中每个节点对应一个 ApiCall 对象。Kernel::ParseGraph的静态签名定义于 codegen_kernel.h。6.2.2 循环与 API 调用生成Loop::Generate按 ImplGraph 的切分策略递归生成循环结构BLOCK_OUT 轴核间并行轴只生成 size/actual_size 计算不生成 for 循环通过 blockidx 隐式并行BLOCK_IN 轴生成block_dim_offset偏移计算确定当前核处理的数据分片其余轴生成for (axis 0; axis loop_size; axis)循环并在循环体内生成尾块tail的实际 size 计算与缓存使能条件。循环体内每遇到一个计算节点按固定骨架生成代码WaitInputs → AllocOutputs → API 调用 → SyncOutputs → FreeInputs → FreeUnusedOutputs即围绕 API 调用完成 UB 队列的等待、分配、同步与释放保证 MTE 搬运与 Vector 计算的正确流水。6.2.3 多模板 tiling_key 分发与传统算子代码类似针对不同的切分策略Schedule 会生成多个模板对应不同的 ImplGraphCodegen 需要依次解析这些模板生成模板函数并在核函数入口处根据 tiling_key 调用不同的模板函数。示例如下extern C __global__ __aicore__ void abs_test(GM_ADDR abs_test_Data_0, GM_ADDR abs_test_Output_0, GM_ADDR workspace, AutofuseTilingData param) { const AutofuseTilingData t; if (t.tiling_key 0) { abs_test_0_general_0_nil_2_nil(abs_test_Data_0, abs_test_Output_0, workspace, t); } else if (t.tiling_key 1) { abs_test_0_general_0_nil_2_nil_unaligned(abs_test_Data_0, abs_test_Output_0, workspace, t); } }上面AutofuseTilingData结构体示例如下其中的成员变量具体值包括 tiling_key 等由 ATT 计算给出BEGIN_TILING_DATA_DEF_T(AutofuseTilingData) const uint32_t block_dim 40; const uint32_t corenum 0; const uint32_t ub_size 196352; const uint32_t hbm_size 0; const uint32_t tiling_key 0; const uint32_t z0z1z2z3t_size 17968; const uint32_t z0z1z2z3Tb_size 57; const uint32_t q0_size 96; const uint32_t q1_size 35936; const uint32_t q2_size 35936; const uint32_t b0_size 35936; const uint32_t tmp_tbuf_size 8192; END_TILING_DATA_DEF_T;不同的模板函数生成的 for 循环以及每次处理的数据量各不相同示例如下tiling_key0对 z1 轴进行切分每次在 UB 中完成计算的向量化轴为 z1t。for (int z0z1Tb 0; z0z1Tb z0z1Tb_loop_size; z0z1Tb) { Abs(y_local[0], xlocal[0], z1t_actual_size); }tiling_key1对 z0 轴进行切分每次在 UB 中完成计算的向量化轴为 z0t 以及 z1。for (int z0Tb 0; z0Tb z0Tb_loop_size; z0Tb) { Abs(y_local[0], xlocal[0], z0t_actual_size * z1_axis_size); }Kernel::GenKernelFuncByTilingKey根据场景选择不同的分发实现对应 codegen_kernel.h 声明寄存器 tiling_key静态 shape将 tiling_key 写入寄存器每个调度模板生成一个 kernel 函数入口处生成if (t.tiling_key N) { ... }分发链当 group 内多模板需要并行时插入SyncAll()同步。其具体实现细分为GenKernelFuncWithRegTilingKey、GenSingleGroupKernelWithRegTilingKey、GenMulGroupKernelWithRegTilingKey三个静态方法解析 tiling_data动态 shape / Inductor运行时解析 tiling_data 结构选择分支对应GenKernelFuncWithParseTilingData及其单/多 group 变体CV 融合生成多 group 的 CV kernel 函数对应GenCVKernelFuncWithMulGroup模板数超限tiling_key 数量超过单编译单元上限时自动拆分为多个编译单元PackingFunc分别编译对应GenPackingFunctions/GetMaxGroupPerCompileUnit等私有方法。6.2.4 API 源码内联Kernel 文件中调用的 AscendC API如DataCopyPadExtend、Abs等的实现源码由AscendCApiRegistryautofuse/codegen/ascendc_api_registry.cpp统一管理。构建期将各 API 头文件打包为源码字符串注册进单例表生成 kernel 时按节点声明的头文件依赖将 API 源码字符串直接内联到 kernel 文件中使每个 kernel 源文件自包含、可独立编译。6.3 Host Code 生成6.3.1 tiling_data 生成TilingData::Generateautofuse/codegen/codegen_tiling_data.cpp生成AutofuseTilingData结构体定义字段包括 tiling_key、block_dim、workspace 大小、各切分轴 size 等多 ScheduleGroup 场景生成统一的包装结构。对于静态 shape 的 Inductor 场景还会生成全 const 初值版本将 tiling 计算前移到编译期。6.3.2 tiling_func 生成Host 侧 tiling 函数由TilingLib组织生成autofuse/codegen/codegen_tiling.cpp包括AutofuseGetTilingSize估算 tiling 所需内存、AutofuseTiling主 tiling 函数根据 shape 计算 tiling_key 和各轴切分大小、workspace 计算等函数。除常规生成外tiling 侧还支持PGO 生成生成独立的 tiling 搜索可执行程序在真实硬件上遍历 tiling 候选、采集性能并保存最优结果Inductor TopN 选择结合性能建模公式由 ATT 提供耗时 ≈ API 调用次数 × 单次耗时与实测 profiling对 tiling 候选做建模评估与 TopN 筛选CV 融合 tiling针对 Cube/Vector 融合场景生成专用 tiling 函数与辅助接口。PGO 相关实现进一步拆分在 autofuse/codegen/codegen_tiling_pgo_search.cpp、codegen_tiling_pgo_memory.cpp、codegen_tiling_pgo_runtime.cpp 等文件中GeneratorPgo入口声明见 codegen.h。6.3.3 infer_shape 生成InfershapeGen::GenInferShapeFuncautofuse/codegen/codegen_infershape.cpp生成extern C InferShape函数解析各输出 shape 表达式中的符号变量生成逐维 AppendDim 的推断代码对含除法取整的表达式额外生成四舍五入与精度校验逻辑保证动态 shape 下形状推断与 kernel 实际计算一致。6.3.4 get_kernel 生成Codegen::GenGetKernelAndJson读取已编译的 kernel 二进制ota生成extern C void GetKernelBin(std::vectorchar)函数将 Device 代码以二进制形式嵌入 Host 侧源码供运行时直接加载免去运行期编译开销。6.4 ApiCall 框架ImplGraph 上每个节点最终都会转化为一次 API 调用这一转换由 ApiCall 框架承载是 Codegen 的核心扩展点。6.4.1 基类与生成链路所有 ApiCall 继承自基类ApiCallautofuse/codegen/codegen_kernel_loop.h核心生成链路为Init(node) // 收集 unit、输入输出 tensor、tmp_buf 信息 → ParseAttr() // 子类覆写解析算子属性 → BuildApiParam() // 组装 CodegenApiParam循环轴、offset、cal_count 等 → GenerateApiCallString()// 按 api_param 拼接 API 调用字符串 → PostProcess() // 生成收尾同步代码其中CodegenApiParam是 API 入参的中间表示将图上的符号表达式ge::Expression经ExpressionConvertStructautofuse/codegen/expression_convert_struct.cpp翻译为 tiling 变量、循环变量的 C 表达式组合例如cal_count、偏移量、循环 size 等。基类还提供 UB 生命周期辅助接口WaitInputs/AllocOutputs/SyncOutputs/FreeInputs等处理共享队列复用与 MTE2/MTE3 同步子类通常只需覆写BuildApiParam和Generate即可接入新算子。6.4.2 工厂与注册机制ApiCall 的创建采用静态自注册工厂模式autofuse/codegen/api_call/utils/api_call_factory.h各 ApiCall 实现类在文件尾定义静态ApiCallRegisterT对象向单例工厂ApiCallFactory注册类名 → 构造函数IR 定义侧ascir codegen impl通过GetApiCallName()声明节点对应的 ApiCall 类名Kernel 生成遍历节点时按类名从工厂创建 ApiCall 对象CreateApiCallObject见 codegen_kernel_loop.cpp主生成流程无需感知具体算子实现新增算子只需增加实现类并注册。6.4.3 内置 ApiCall主框架autofuse/codegen/api_call/内置的 ApiCall 按计算类别组织类别ApiCall数据搬运LoadApiCall、StoreApiCall一元计算UnaryApiCall、CastApiCall、RsqrtApiCall、LogicalNotApiCall、UnaryBitWidthChangeApiCall等二元计算BinaryApiCall、AxpyApiCall、TrueDivApiCall、PowApiCall、LeakyReluApiCall等比较与选择CompareApiCall、WhereApiCall、ClipByValueApiCall归约ReduceApiCall数据重排TransposeApiCall、BroadcastApiCall、ConcatApiCall、GatherApiCall其他PadApiCall、RemovePadApiCall、Ub2UbApiCall等主框架内置实现的目录组织为 autofuse/codegen/api_call/ 下的 broadcast、concat、datacopy、elewise、gather、reduce、transpose、utils 等子目录与上表类别一一对应。6.5 v35 平台扩展层不同 SoC 版本的硬件指令集和 API 形态存在差异Codegen 通过平台扩展层autofuse/v35/codegen/承载版本特定实现。扩展层完全复用主框架的基类与注册机制各实现类同样继承ApiCall并静态注册进ApiCallFactory由 SoC 版本路由在 IR 注册表层面按芯片版本选择实现Kernel 生成流程本身保持平台无关。v35 扩展层包含四类实现6.5.1 reg_api_callRegBase 风格 API面向 RegTensor 指令形态的向量计算与数据搬运 API是 v35 的主体包括BinaryApiCallV2、CastV2ApiCall、CompareV2ApiCall、LoadRegApiCall、StoreRegApiCall、NddmaApiCall、IndirectLoadRegApiCall、BroadcastRegApiCall、ConcatRegApiCall、SplitRegApiCall、RegReduceApiCall、TransposeRegApiCall、GatherRegApiCall、WhereRegApiCall、SoftmaxApiCall等约 24 个类。相比主框架版本V2 系列 API 具备更强的参数化能力如轴合并 merge_axes并会在BuildApiParam中计算outer_call_countAPI 外层调用次数该值直接对接 Auto Tiling 的性能公式建模耗时 ≈ 调用次数 × 单次耗时作为 tiling 候选评估的输入。对应的 API 源码模板约 100 个 RegBase API由构建期从源码头文件自动打包为字符串注册进AscendCApiRegistryautofuse/v35/codegen/ascendc_reg_base_api_register.cpp。6.5.2 cube_api_callCube 算子 APIMatmulApiCall、Conv2DApiCall生成 Matmul/Conv 等 Cube 算子的 AscendC API 调用。Cube 单元调用除生成调用点外还会生成模板函数定义GenerateFuncDefinition和宏展开GenerateMacro并通过 tiling wrapper 对接官方 matmul tiling 实现。相关 Cube API 模板matmul、batch_matmul、conv2d 及其 dynamic/tiling_key/pingpong 变体注册见 autofuse/v35/codegen/ascendc_cube_api_register.cpp。6.5.3 micro_api_callMicroAPIMicroAPI 是最低一级的代码形态直接生成AscendC::MicroAPI::*指令级调用如LoadAlign、StoreAlign与AscendC::Reg::*寄存器操作如Pack、UnPack、MaskPack由独立的MicroApiCall基类与MicroApiCallFactory工厂承载不继承主框架ApiCall包括MicroLoadApiCall、MicroStoreApiCall、MicroCastApiCall、MicroCompareApiCall、MicroWhereApiCall、MicroBinaryScalarApiCall、MicroScalarBroadcastApiCall等。6.5.4 vec_func_callVectorFuncVfCall将 ImplGraph 中的子图封装为一个独立的 vf 函数实现多次循环 → 一次函数调用的结构优化。其内部持有VFLoop循环树遍历子图节点时通过 MicroAPI 工厂为每个节点创建 MicroApiCall形成两层结构VfCall主框架 ApiCall生成函数定义与调用点 └── VFLoopvf 内循环树 └── MicroApiCall指令级 API 调用该机制同时支持 CV-UB 融合场景的 kernel 生成与Kernel::GenerateVecFuncOfCVFusion、InitCVFusionAddr等 CV 融合相关实现codegen_kernel.h配套工作。6.6 辅助机制API 参数 dumpCodegenApiParam::DumpGraphApiParamsautofuse/codegen/codegen_api_param_dump.cpp可将每张 ImplGraph 中所有节点的 API 入参序列化为文本文件用于调试和对比生成的调用参数开发新 ApiCall 时可通过该工具核验参数组装是否正确表达式转换ExpressionItem/CombinedExpression提供符号表达式到 C 代码的统一渲染并附带工厂函数ExprItemFactory/CombinedExprFactory用于组装常见表达式模式如SizeMinusActualSize尾块计算。7. 关键文件索引文件路径职责autofuse/codegen/codegen.h / codegen.cppCodegen 门面类总入口与产物组装autofuse/codegen/codegen_kernel.h / codegen_kernel.cppKernel 生成器、Tiler/TPipe 数据模型、tiling_key 分发autofuse/codegen/codegen_kernel_loop.h / codegen_kernel_loop.cppLoop 循环树、ApiCall 基类与生成骨架autofuse/codegen/codegen_tiling.cpp 等Host 侧 tiling 生成含 PGO、Inductor TopN、CV tilingautofuse/codegen/codegen_tiling_data.cppAutofuseTilingData 结构体生成autofuse/codegen/codegen_infershape.cppInferShape 函数生成autofuse/codegen/codegen_graph_check.cppImplGraph 合法性校验autofuse/codegen/expression_convert_struct.cpp符号表达式到 C 表达式转换autofuse/codegen/codegen_api_param_dump.cppAPI 入参 dump 调试工具autofuse/codegen/ascendc_api_registry.cppAscendC API 源码模板注册表autofuse/codegen/api_call/主框架内置 ApiCall 实现datacopy/elewise/reduce/transpose/gather/broadcast/concatautofuse/v35/codegen/reg_api_call/RegBase 风格 V2 ApiCall平台扩展层主体autofuse/v35/codegen/cube_api_call/Matmul/Conv2D Cube ApiCallautofuse/v35/codegen/micro_api_call/MicroAPI 指令级 ApiCall 及独立工厂autofuse/v35/codegen/vec_func_call/VectorFunc 子图封装VfCall/VFLoopautofuse/v35/codegen/ascendc_reg_base_api_register.cppRegBase API 模板注册autofuse/v35/codegen/ascendc_cube_api_register.cppCube API 模板注册8. 小结Codegen 以生成代码的形、不决定 tiling 的值为根本设计原则通过 ImplGraph 这一唯一接口与 Schedule 解耦通过AutofuseTilingData与 ATT 分工通过 ApiCall 静态注册工厂实现算子接入的平台无关扩展再以 v35 平台扩展层承载不同 SoC 版本的指令形态差异。理解这一套契约 门面 扩展点的架构是向 graph-autofusion 中新增融合算子、适配新 SoC 版本或扩展 tiling 策略的起点而 Codegen 设计文档 本身则是最佳的架构导览本文所列源码路径可作为逐模块深入阅读的索引。【免费下载链接】graph-autofusionGraph-autofusion 是一个面向昇腾Ascend芯片的轻量级、解耦式组件集合旨在通过自动融合技术加速模型执行。 目前已开源 SuperKernel 组件和 Autofuse 组件未来将持续开放更多自动融合相关模块。项目地址: https://gitcode.com/cann/graph-autofusion创作声明：本文部分内容由AI辅助生成（AIGC），仅供参考