CANN Runtime 实战用 aclrtBinaryEnumerateFunctions 在单个算子二进制中枚举并启动多个 Kernel【免费下载链接】runtime本项目提供CANN运行时组件和维测功能组件。项目地址: https://gitcode.com/cann/runtime本篇技术指南围绕 CANN Runtime 开源仓库中的binary_enumerate_functions示例展开讲解如何将多个 Device Kerneladd_custom、sub_custom、mul_custom编译进同一个独立算子二进制文件并通过aclrtBinaryEnumerateFunctions一次性枚举出全部函数句柄、用aclrtGetFunctionName反查函数名再逐一启动执行。读完本文你将掌握多 Kernel 单二进制加载与启动的完整调用链并能在自己的算子工程中复刻这套一包多函数的交付与运行模式。一、示例要解决的核心问题常规的 Kernel 启动流程如0_launch_kernel示例通常将单个 Kernel 编译为独立二进制Host 侧加载后拿到函数句柄再下发。当一个工程内含多个算子例如向量加、减、乘时逐个编译、逐个加载会带来多个二进制文件的管理与传输开销。2_binary_enumerate_functions示例展示了另一种更紧凑的交付形态把多个 Kernel 函数编译进同一个独立算子二进制Host 侧只加载这一个文件然后通过aclrtBinaryEnumerateFunctions枚举出模块内全部函数句柄循环启动它们。该示例目录结构如下见 README.md2_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 侧使用x1.0、y2.0作为输入最终三个 Kernel 分别输出3.0、-1.0和2.0正好对应加、减、乘三种运算。二、Kernel 侧实现一份代码三个导出函数Device Kernel 源码位于 kernel/custom.cpp采用模板化的方式避免三份重复代码使用enum class BinaryOperation { ADD, SUB, MUL }标记运算类型模板类KernelBinaryoperation内部基于 AscendC 的TPipe、TQueQuePosition::VECIN/VECOUT、GlobalTensorhalf完成标准的 CopyIn → Compute → CopyOut 流水Init中通过AscendC::GetBlockIdx()按 block 划分全局内存区间BLOCK_LENGTH 8 * 2048 / 8并初始化输入/输出队列BUFFER_NUM 2双缓冲TILE_LENGTH BLOCK_LENGTH / TILE_NUM / BUFFER_NUMTILE_NUM 8Compute中使用if constexpr按模板参数选择AscendC::Add/AscendC::Sub/AscendC::Mul完成向量运算。文件末尾导出三个extern C __global__ __aicore__入口函数add_custom、sub_custom、mul_custom每个函数实例化对应运算的KernelBinary并调用InitProcess。这三个函数会被 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为nullptrbinHandle作为输出接收二进制句柄。加载阶段只在 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 准备输入数据并启动三个 Kernelstd::vectoraclFloat16 x(kElementCount, aclFloatToFloat16(1.0F)); std::vectoraclFloat16 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])); }启动采用aclrtLaunchKernelWithHostArgsacl_rt.h#L5270-L5272以参数数组形式直接把三个设备地址打包传给 KernelaclError 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保证即便某一步失败也能继续清理其余资源。四、产品支持情况示例 READMEREADME_en.md明确列出支持的产品产品是否支持Ascend 950PR / Ascend 950DT√Atlas A3 训练系列产品 / Atlas A3 推理系列产品√Atlas A2 训练系列产品 / Atlas A2 推理系列产品√五、环境准备与一键构建运行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 环境优先使用预设变量ASCEND_INSTALL_PATH/ASCEND_HOME_PATH两者互为兜底run.sh#L28-L29若未预设则 source example/common/resolve_cann_env.sh 在常见安装路径中查找并 sourceset_env.sh当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_VERSIONAscend910B15.4 编译与运行cd ${git_clone_path}/example/2_advanced_features/kernel/2_binary_enumerate_functions bash run.shrun.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]: nameadd_custom, handle... [INFO] function[1]: namemul_custom, handle... [INFO] function[2]: namesub_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创作声明：本文部分内容由AI辅助生成（AIGC），仅供参考 CANN Runtime 实战：用 aclrtBinaryEnumerateFunctions 在单个算子二进制中枚举并启动多个 Kernel
CANN Runtime 实战用 aclrtBinaryEnumerateFunctions 在单个算子二进制中枚举并启动多个 Kernel【免费下载链接】runtime本项目提供CANN运行时组件和维测功能组件。项目地址: https://gitcode.com/cann/runtime本篇技术指南围绕 CANN Runtime 开源仓库中的binary_enumerate_functions示例展开讲解如何将多个 Device Kerneladd_custom、sub_custom、mul_custom编译进同一个独立算子二进制文件并通过aclrtBinaryEnumerateFunctions一次性枚举出全部函数句柄、用aclrtGetFunctionName反查函数名再逐一启动执行。读完本文你将掌握多 Kernel 单二进制加载与启动的完整调用链并能在自己的算子工程中复刻这套一包多函数的交付与运行模式。一、示例要解决的核心问题常规的 Kernel 启动流程如0_launch_kernel示例通常将单个 Kernel 编译为独立二进制Host 侧加载后拿到函数句柄再下发。当一个工程内含多个算子例如向量加、减、乘时逐个编译、逐个加载会带来多个二进制文件的管理与传输开销。2_binary_enumerate_functions示例展示了另一种更紧凑的交付形态把多个 Kernel 函数编译进同一个独立算子二进制Host 侧只加载这一个文件然后通过aclrtBinaryEnumerateFunctions枚举出模块内全部函数句柄循环启动它们。该示例目录结构如下见 README.md2_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 侧使用x1.0、y2.0作为输入最终三个 Kernel 分别输出3.0、-1.0和2.0正好对应加、减、乘三种运算。二、Kernel 侧实现一份代码三个导出函数Device Kernel 源码位于 kernel/custom.cpp采用模板化的方式避免三份重复代码使用enum class BinaryOperation { ADD, SUB, MUL }标记运算类型模板类KernelBinaryoperation内部基于 AscendC 的TPipe、TQueQuePosition::VECIN/VECOUT、GlobalTensorhalf完成标准的 CopyIn → Compute → CopyOut 流水Init中通过AscendC::GetBlockIdx()按 block 划分全局内存区间BLOCK_LENGTH 8 * 2048 / 8并初始化输入/输出队列BUFFER_NUM 2双缓冲TILE_LENGTH BLOCK_LENGTH / TILE_NUM / BUFFER_NUMTILE_NUM 8Compute中使用if constexpr按模板参数选择AscendC::Add/AscendC::Sub/AscendC::Mul完成向量运算。文件末尾导出三个extern C __global__ __aicore__入口函数add_custom、sub_custom、mul_custom每个函数实例化对应运算的KernelBinary并调用InitProcess。这三个函数会被 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为nullptrbinHandle作为输出接收二进制句柄。加载阶段只在 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 准备输入数据并启动三个 Kernelstd::vectoraclFloat16 x(kElementCount, aclFloatToFloat16(1.0F)); std::vectoraclFloat16 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])); }启动采用aclrtLaunchKernelWithHostArgsacl_rt.h#L5270-L5272以参数数组形式直接把三个设备地址打包传给 KernelaclError 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保证即便某一步失败也能继续清理其余资源。四、产品支持情况示例 READMEREADME_en.md明确列出支持的产品产品是否支持Ascend 950PR / Ascend 950DT√Atlas A3 训练系列产品 / Atlas A3 推理系列产品√Atlas A2 训练系列产品 / Atlas A2 推理系列产品√五、环境准备与一键构建运行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 环境优先使用预设变量ASCEND_INSTALL_PATH/ASCEND_HOME_PATH两者互为兜底run.sh#L28-L29若未预设则 source example/common/resolve_cann_env.sh 在常见安装路径中查找并 sourceset_env.sh当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_VERSIONAscend910B15.4 编译与运行cd ${git_clone_path}/example/2_advanced_features/kernel/2_binary_enumerate_functions bash run.shrun.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]: nameadd_custom, handle... [INFO] function[1]: namemul_custom, handle... [INFO] function[2]: namesub_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创作声明：本文部分内容由AI辅助生成（AIGC），仅供参考