运行时加载机制

本文是扩展内容，介绍算子开发完成后，编译产物如何被加载和执行，以及开发时的设计决策如何影响运行表现。本文从算子开发者视角梳理运行时调用的基本模型和代码写法对运行行为的影响。

1. 算子是如何被调用的

二段式接口

单算子API采用二段式调用：aclnnXxxGetWorkspaceSize和aclnnXxx。

第一阶段：GetWorkspaceSize

C++

aclnnXxxGetWorkspaceSize(input0, input1, output, &workspaceSize, &executor);

此阶段主要在Host侧完成执行前准备：

1. 根据OpType、SoC、输入输出dtype/format/shape、属性等信息，优先尝试匹配aclnn参数缓存；缓存命中时复用已保存的binary和Tiling执行信息，跳过后续binary匹配和TilingFunc调用。
2. 缓存未命中时，根据算子注册信息和当前参数匹配可用的核函数（Kernel）binary，或者调用算子注册的TilingFunc计算当前输入对应的TilingData、TilingKey、BlockDim。
3. 根据匹配到的binary配置和TilingKey，确定本次执行的核函数（Kernel）入口。
4. 返回workspaceSize（需要分配的Device内存大小）和executor（承载调度上下文）。

此阶段可能触发的耗时操作：

• TilingFunc调用：执行Host侧C++函数。
• Binary查找：在文件系统中检索.o文件，涉及文件I/O。

第二阶段：Execute

C++

aclnnXxx(workspace, workspaceSize, executor, stream);

此阶段将第一阶段准备好的全部信息打包提交到stream上执行。

为什么分两阶段

两阶段接口的目的是将执行前准备和执行下发分离：调用方可以先获取workspaceSize和executor，再根据自身资源管理方式选择按算子申请、统一内存池分配或复用已有workspace，并在提交核函数（Kernel）前完成内存预算，可以减少频繁申请内存的开销和减少内存碎片。

Binary匹配机制

运行时框架会基于OpType、目标芯片型号、输入输出dtype/format、shape支持范围、属性以及TilingKey/模板参数组合等信息匹配预编译的核函数（Kernel）binary。

binary查找的优先级顺序：

• 运行时会先加载ASCEND_CUSTOM_OPP_PATH配置路径下的算子包，再查找默认CANN路径下的算子包。
• 默认安装到$ASCEND_OPP_PATH/vendors/<vendor_name>的场景下，算子包优先级根据opp/vendors/config.ini中的load_priority决定，排在前面的算子包优先级更高。
• 动态库/静态库集成场景的加载方式与RUN包不同，应用直接链接或通过ASCEND_CUSTOM_OPP_PATH配置动态库时，应按算子动态库和静态库编译中的优先级规则处理。

2. 代码写法如何影响运行行为

开发者在编写算子代码时做出的选择，会直接影响编译产物的形态和运行时的表现。以下按影响阶段分为三类。

影响编译产物形态的写法

TilingKey / 核函数（Kernel）模板的分支数量

算子有两种分支策略可选：数字TilingKey或命名核函数（Kernel）模板参数，两者互斥。

分支策略	编译耗时	运行表现	适用场景
单一分支（所有支持输入走同一套核函数（Kernel）逻辑）	极快	代码简单，但难以针对不同场景优化。	简单算子、开发阶段。
按shape分档（如 ≤64 / ≤1024 / ≤64K / >64K）	适中	可针对不同数据规模选择更合适的算法或tile参数。	大多数算子。
按shape × dtype × 参数组合精细分档	较长	性能特化更充分，但binary数量和包体积增加。	生产部署、性能敏感。

开发阶段建议：少分支或用选择性编译选项只编译当前验证的分支，加速迭代。生产部署：确保已发布的二进制配置覆盖实际会调用到的dtype/format、属性、TilingKey或模板参数组合。

ASCEND_SOC_SERIES/ASCEND_COMPUTE_UNIT范围

离线编译产物需要包含当前运行芯片型号对应的核函数（Kernel）binary。开发阶段如果只验证单一芯片，可收窄ASCEND_SOC_SERIES或ASCEND_COMPUTE_UNIT；两者不能同时配置。需要跨芯片部署时，再按目标环境配置多个系列或型号。ASCEND_SOC_SERIES的输入大小写不敏感；产物目录和运行时查找规则不变。如果离线编译产物里不包含当前运行芯片的核函数（Kernel）binary，运行时会失败。

影响Host侧运行开销的写法

TilingFunc的计算复杂度

TilingFunc被调用时，如果函数中包含复杂计算（如大量循环、动态内存分配、文件系统操作），会直接增加每次算子调用的Host侧延迟。

建议：TilingFunc中只做必要的标量计算（shape推导、分块参数计算），避免复杂逻辑。

TilingData字段数量

TilingData结构体在Host侧填写、序列化后传递到核函数（Kernel）侧。字段越多，序列化开销越大。

建议：TilingData只传核函数（Kernel）需要的标量值，不要传冗余信息。

Workspace大小设置

context->GetWorkspaceSizes()设置的workspace大小决定了每次调用需要额外分配的Device内存量。workspace越大，内存压力越大。

建议：精确计算所需workspace大小，避免预留过多冗余空间。

影响核函数（Kernel）侧执行表现的写法

TILING_KEY_IS vs运行时if-else

C++

// 方式一：TilingKey分支（推荐）
if (TILING_KEY_IS(1)) {
    ProcessSmall();    // 编译期常量折叠，未命中的分支被完全移除
} else if (TILING_KEY_IS(2)) {
    ProcessLarge();
}

// 方式二：运行时分支
if (tilingData.totalLength <= threshold) {
    ProcessSmall();    // 运行时条件判断，两个分支都在binary中
} else {
    ProcessLarge();
}

TILING_KEY_IS是编译期常量判断，编译器将未命中的分支完全移除（常量折叠），消除运行时分支判断开销和icache压力。运行时if-else两个分支都保留在binary中，代码体积更大。

核函数（Kernel）模板参数

核函数（Kernel）模板使用命名参数（如D_T_X、TILE_NUM）组织模板参数组合，并在Host侧根据实际参数生成和设置TilingKey，从而实现编译期特化：

C++

template<typename D_T_X, typename D_T_Y, typename D_T_Z, int TILE_NUM>
__global__ __aicore__ void add_custom(...)
{
    // D_T_X/D_T_Y/D_T_Z在编译期确定，类型判断在编译期完成
    // TILE_NUM在编译期确定，循环边界等可直接展开
}

模板参数使编译器能够针对特定类型和常量做优化（如循环展开、常量传播）。运行时只需选择已编译的模板组合，不会在核函数（Kernel）内产生模板参数本身的判断开销。

tileNum切分粒度

TilingData中的tileNum决定每个Block上的分块个数：

• tileNum过小 → 每个分块的数据量大，单次搬运耗时长，但循环次数少。
• tileNum过大 → 每个分块的数据量小，循环次数多，循环开销占比增大。

需要根据Unified Buffer（UB）内存容量和计算/搬运比例选取合适的tileNum。

相关文档

• 基本流程 — 编译与部署操作指南。
• Host侧Tiling实现 — Tiling函数、TilingData和workspace设置方法。
• 多分支策略 — TilingKey和核函数（Kernel）模板的配置方法。
• 单算子API调用 — 两段式接口的使用方法。
• 基于simplified key的运行时选择流程 — 回调生成simplified key并选择预编译binary的流程。
• 算子动态库和静态库编译 — 动态库/静态库形态下的集成方式。