企业官网建设流程全解析

CANN Runtime 实战：用 aclrtBinaryEnumerateFunctions 在单个算子二进制中枚举并启动多个 Kernel

【免费下载链接】runtime本项目提供CANN运行时组件和维测功能组件。项目地址: https://gitcode.com/cann/runtime

本篇技术指南围绕 CANN Runtime 开源仓库中的binary_enumerate_functions示例展开，讲解如何将多个 Device Kernel（add_custom、sub_custom、mul_custom）编译进同一个独立算子二进制文件，并通过aclrtBinaryEnumerateFunctions一次性枚举出全部函数句柄、用aclrtGetFunctionName反查函数名，再逐一启动执行。读完本文，你将掌握多 Kernel 单二进制加载与启动的完整调用链，并能在自己的算子工程中复刻这套"一包多函数"的交付与运行模式。

一、示例要解决的核心问题

常规的 Kernel 启动流程（如0_launch_kernel示例）通常将单个 Kernel 编译为独立二进制，Host 侧加载后拿到函数句柄再下发。当一个工程内含多个算子（例如向量加、减、乘）时，逐个编译、逐个加载会带来多个二进制文件的管理与传输开销。

2_binary_enumerate_functions示例展示了另一种更紧凑的交付形态：把多个 Kernel 函数编译进同一个独立算子二进制，Host 侧只加载这一个文件，然后通过aclrtBinaryEnumerateFunctions枚举出模块内全部函数句柄，循环启动它们。

该示例目录结构如下（见 README.md）：

2_binary_enumerate_functions/ ├── CMakeLists.txt ├── kernel/ │ └── custom.cpp # 三个 Device Kernel 的实现 ├── main.cpp # Host 侧加载、枚举、启动 ├── README.md ├── README_en.md └── run.sh # 环境探测 + 构建 + 运行一键脚本

三个 Device Kernel 全部定义在kernel/custom.cpp中；Host 侧使用x=1.0、y=2.0作为输入，最终三个 Kernel 分别输出3.0、-1.0和2.0，正好对应加、减、乘三种运算。

二、Kernel 侧实现：一份代码，三个导出函数

Device Kernel 源码位于 kernel/custom.cpp，采用模板化的方式避免三份重复代码：

• 使用enum class BinaryOperation { ADD, SUB, MUL }标记运算类型；
• 模板类KernelBinary<operation>内部基于 AscendC 的TPipe、TQue<QuePosition::VECIN/VECOUT>、GlobalTensor<half>完成标准的 CopyIn → Compute → CopyOut 流水：
  • Init中通过AscendC::GetBlockIdx()按 block 划分全局内存区间（BLOCK_LENGTH = 8 * 2048 / 8），并初始化输入/输出队列（BUFFER_NUM = 2双缓冲，TILE_LENGTH = BLOCK_LENGTH / TILE_NUM / BUFFER_NUM，TILE_NUM = 8）；
  • Compute中使用if constexpr按模板参数选择AscendC::Add/AscendC::Sub/AscendC::Mul完成向量运算。
• 文件末尾导出三个extern "C" __global__ __aicore__入口函数：add_custom、sub_custom、mul_custom，每个函数实例化对应运算的KernelBinary并调用Init+Process。

这三个函数会被 CMakeLists.txt 中的ascendc_fatbin_library(custom_kernels kernel/custom.cpp)一并编译到名为custom_kernels的独立算子二进制中，编译产物位于out/fatbin/custom_kernels/custom_kernels.o。

三、Host 侧调用链：加载 → 枚举 → 反查 → 启动

Host 逻辑集中在 main.cpp 的RunSample函数（main.cpp#L35-L72），关键步骤依次为：

3.1 初始化 Runtime 与加载二进制

CHECK_ERROR(aclInit(nullptr)); CHECK_ERROR(aclrtSetDevice(kDeviceId)); // kDeviceId = 0 CHECK_ERROR(aclrtCreateStream(&resources.stream)); CHECK_ERROR(aclrtBinaryLoadFromFile(binaryPath, nullptr, &resources.binHandle));

aclrtBinaryLoadFromFile完成 Host 侧的文件加载与解析。该 API 在 acl_rt.h 中声明为：

aclError aclrtBinaryLoadFromFile( const char* binPath, aclrtBinaryLoadOptions* options, aclrtBinHandle* binHandle);

本示例传入的options为nullptr，binHandle作为输出接收二进制句柄。加载阶段只在 Host 侧解析，尚未把数据拷贝到 Device。

3.2 枚举函数句柄并反查函数名

aclrtFuncHandle functions[kFunctionCount] = {}; // kFunctionCount = 3 char functionNames[kFunctionCount][kFunctionNameLength] = {}; // 每项 128 字节 CHECK_ERROR(aclrtBinaryEnumerateFunctions(resources.binHandle, functions, kFunctionCount)); for (uint32_t i = 0U; i < kFunctionCount; ++i) { CHECK_ERROR(aclrtGetFunctionName(functions[i], kFunctionNameLength, functionNames[i])); INFO_LOG("function[%u]: name=%s, handle=%p", i, functionNames[i], functions[i]); }

aclrtBinaryEnumerateFunctions在 acl_rt.h 中的签名为：

aclError aclrtBinaryEnumerateFunctions( aclrtBinHandle const binHandle, aclrtFuncHandle* funcHandles, uint32_t numFunctions);

根据头文件注释（acl_rt.h#L3576-L3589），该 API 用于枚举二进制模块中的函数句柄，且首次访问该二进制句柄时，Runtime 会把关联的算子二进制数据拷贝到当前 Context 对应的 Device——这正是文档中"第一次aclrtBinaryEnumerateFunctions调用完成二进制上载"这一行为的源码级依据。输出数组会被填充min(numFunctions, 实际函数数)个条目，因此调用者传入的数组大小（此处为 3）要与模块内函数数匹配。

随后用aclrtGetFunctionName(funcHandle, maxLen, name)（acl_rt.h#L4355）按句柄反查 Kernel 名，maxLen传入 128 表示缓冲区长度。由于二进制内的函数顺序由编译布局决定，枚举得到的顺序与源码顺序不一定一致，因此"按名反查"是区分各句柄用途的必要手段。

3.3 准备输入数据并启动三个 Kernel

std::vector<aclFloat16> x(kElementCount, aclFloatToFloat16(1.0F)); std::vector<aclFloat16> y(kElementCount, aclFloatToFloat16(2.0F)); const size_t dataSize = kElementCount * sizeof(aclFloat16); // 8 * 2048 个 half 元素 CHECK_ERROR(aclrtMalloc(&resources.xDevice, dataSize, ACL_MEM_MALLOC_HUGE_FIRST)); // ... yDevice、zDevice 同理，并 aclrtMemcpy 完成 H2D 拷贝 for (uint32_t i = 0U; i < kFunctionCount; ++i) { void* args[] = {resources.xDevice, resources.yDevice, resources.zDevice}; CHECK_ERROR(aclrtLaunchKernelWithHostArgs( functions[i], kBlockDim, resources.stream, nullptr, args, sizeof(args), nullptr, 0U)); CHECK_ERROR(aclrtSynchronizeStream(resources.stream)); CHECK_ERROR(aclrtMemcpy(z.data(), dataSize, resources.zDevice, dataSize, ACL_MEMCPY_DEVICE_TO_HOST)); INFO_LOG("%s result: %.1f", functionNames[i], aclFloat16ToFloat(z[0])); }

启动采用aclrtLaunchKernelWithHostArgs（acl_rt.h#L5270-L5272），以"参数数组"形式直接把三个设备地址打包传给 Kernel：

aclError aclrtLaunchKernelWithHostArgs( aclrtFuncHandle funcHandle, uint32_t numBlocks, aclrtStream stream, aclrtLaunchKernelCfg* cfg, void* hostArgs, size_t argsSize, aclrtPlaceHolderInfo* placeHolderArray, size_t placeHolderNum);

其中numBlocks为kBlockDim = 8（与 Kernel 内USE_CORE_NUM = 8对应，8 个 AI Vector 核并行切分数据），cfg与placeHolderArray均传nullptr/0。每次启动后紧跟aclrtSynchronizeStream同步，再把结果 D2H 拷回 Host，打印首个元素z[0]作为正确性验证。

3.4 资源释放

ReleaseResources按逆序释放全部资源：aclrtFree释放三块设备内存、aclrtBinaryUnLoad卸载二进制、aclrtDestroyStreamForce销毁流、aclrtResetDeviceForce复位设备、aclFinalize反初始化。释放阶段统一使用CHECK_ERROR_WITHOUT_RETURN（见 utils.h），保证即便某一步失败也能继续清理其余资源。

四、产品支持情况

示例 README（README_en.md）明确列出支持的产品：

五、环境准备与一键构建运行

5.1 环境要求

• CANN Runtime 默认安装路径为/home/developer/Ascend/cann；
• 构建 Device Kernel 还需要AscendC 编译工具。仅安装 Runtime 包时可能不存在ascendc.cmake，此时需额外安装并 source CANN Toolkit 环境（run.sh 中会显式检查${ASCENDC_CMAKE_DIR}/ascendc.cmake是否存在，缺失即报错退出）。

5.2 环境自动探测机制

run.sh会按以下优先级解析 CANN 环境：

1. 优先使用预设变量ASCEND_INSTALL_PATH/ASCEND_HOME_PATH（两者互为兜底，run.sh#L28-L29）；
2. 若未预设，则 source example/common/resolve_cann_env.sh 在常见安装路径中查找并 sourceset_env.sh；
3. 当SOC_VERSION与ASCENDC_CMAKE_DIR均未设置时，source example/set_sample_env.sh 自动探测：该脚本会临时编译 example/tools/get_soc_version/get_soc_version.cpp 作为辅助程序，通过 ACL API 从设备查询 SOC 版本号，并按宿主架构（x86_64-linux/aarch64-linux）在 CANN 包布局中定位包含ascendc.cmake的tikcpp/ascendc_kernel_cmake目录，最后导出ASCEND_INSTALL_PATH、ASCEND_HOME_PATH、SOC_VERSION、ASCENDC_CMAKE_DIR四个变量。

因此，在设备可达、CANN 安装完整的机器上，无需任何手工配置即可直接运行。

5.3 手动配置方式（可选）

export ASCEND_INSTALL_PATH=/home/developer/Ascend/cann export SOC_VERSION=Ascend910B1

5.4 编译与运行

cd ${git_clone_path}/example/2_advanced_features/kernel/2_binary_enumerate_functions bash run.sh

run.sh内部依次完成：清理旧的build/与out/→ 调用cmake -S . -B build→cmake --build→cmake --install→ 校验 Kernel 二进制与 Host 可执行文件是否生成 → 按宿主架构设置LD_LIBRARY_PATH（含${ASCEND_INSTALL_PATH}/runtime/lib64、${ASCEND_INSTALL_PATH}/lib64、${ASCEND_INSTALL_PATH}/${arch_dir}/lib64）→ 执行binary_enumerate_functions custom_kernels.o。

5.5 预期输出

成功运行时输出类似（README 原样给出，handle地址因环境而异）：

[INFO] Enumerating functions in the Kernel binary. [INFO] aclrtBinaryEnumerateFunctions succeeded. [INFO] function[0]: name=add_custom, handle=... [INFO] function[1]: name=mul_custom, handle=... [INFO] function[2]: name=sub_custom, handle=... [INFO] add_custom result: 3.0 [INFO] mul_custom result: 2.0 [INFO] sub_custom result: -1.0

输出解读：

• 三个函数句柄枚举成功，函数名与kernel/custom.cpp中的导出符号一一对应（注意枚举顺序是add → mul → sub，与源码声明顺序不同，印证了"必须按名反查"的必要性）；
• 结果验证：x + y = 1.0 + 2.0 = 3.0、x * y = 1.0 * 2.0 = 2.0、x - y = 1.0 - 2.0 = -1.0，三个 Kernel 均正确执行。

六、扩展阅读与关联示例

• 二进制函数计数：aclrtBinaryGetFunctionCount(binHandle, &count)（acl_rt.h#L3560-L3561）可在枚举前获取模块内函数总数，配套示例见 3_binary_get_function_count；
• 按入口地址取句柄：aclrtBinaryGetFunctionByEntry(binHandle, funcEntry, &funcHandle)（acl_rt.h#L3549-L3550）支持以函数入口寻址；
• 符号反查：aclrtGetFuncBySymbol(symbol, &funcHandle)（acl_rt.h#L3599）提供按符号定位句柄的另一条路径；
• 基础启动对比：单 Kernel 加载与启动流程可参考 0_launch_kernel，通过对照可更清晰理解"多函数单二进制"模式相对逐文件加载的差异。

七、小结

通过本示例可以看到，CANN Runtime 的"算子二进制"是一个可容纳多个 Kernel 函数的模块化载体：aclrtBinaryLoadFromFile负责 Host 侧解析，aclrtBinaryEnumerateFunctions首次调用即触发二进制上载并枚举全部函数句柄，aclrtGetFunctionName用于按名识别，最终由aclrtLaunchKernelWithHostArgs统一调度执行。这种模式既减少了二进制文件数量，也让 Host 侧代码可以用统一循环驱动多个算子，适用于算子库、多形态融合 Kernel 等"一包多函数"的交付场景。

【免费下载链接】runtime本项目提供CANN运行时组件和维测功能组件。项目地址: https://gitcode.com/cann/runtime