CANN Runtime 算子二进制核函数数量查询aclrtBinaryGetFunctionCount 实战指南【免费下载链接】runtime本项目提供CANN运行时组件和维测功能组件。项目地址: https://gitcode.com/cann/runtimeaclrtBinaryGetFunctionCount是 CANN Runtime 提供的用于查询算子二进制Binary中核函数Kernel Function总数的接口。本文以开源仓库 3_binary_get_function_count 样例 为完整骨架结合其源码实现与头文件声明讲解如何将多个 AscendC 核函数编译进同一个独立算子二进制查询核函数总数后按名称获取句柄、查询函数信息、下发任务并校验结果帮助读者掌握二进制级的 Kernel 管理与调用全流程。样例概述一个二进制文件承载三个核函数在常规开发中一个算子往往对应一个独立编译产物。但当多个核函数如add_custom、sub_custom、mul_custom需要共享数据、复用内存布局或减少加载开销时将它们编译到同一个独立算子二进制文件中再通过统一句柄按名称调度是更高效的组织方式。本样例正是围绕这一场景设计的样例将add_custom、sub_custom和mul_custom三个 AscendC 核函数编译到同一个独立算子二进制文件中随后在 Host 侧完成如下完整链路使用aclrtBinaryLoadFromFile加载算子二进制文件并获得binHandle使用aclrtBinaryGetFunctionCount查询其中核函数总数按核函数名称调用aclrtBinaryGetFunction获取对应句柄调用aclrtGetFunctionName与aclrtGetFunctionAddr查询函数名与 Device 侧地址准备 FP16 输入数据通过aclrtLaunchKernelWithHostArgs分别下发三个核函数逐项校验计算结果最后卸载二进制并释放所有资源。其主程序入口位于 main.cpp三个核函数实现位于 kernel/custom.cpp。产品支持情况本样例支持以下产品来自样例 README 的产品支持表产品是否支持Ascend 950PR/Ascend 950DT√Atlas A3 训练系列产品/Atlas A3 推理系列产品√Atlas A2 训练系列产品/Atlas A2 推理系列产品√目录结构与构建方式样例目录结构如下3_binary_get_function_count/ ├── CMakeLists.txt # 编译脚本同时构建 AscendC 算子二进制与 Host 程序 ├── kernel/ │ └── custom.cpp # 三个 AscendC 核函数add_custom/sub_custom/mul_custom ├── main.cpp # Host 侧主程序展示完整调用流程 ├── README.md ├── README_en.md └── run.sh # 一键编译运行脚本含环境自动探测在 CMakeLists.txt 中算子二进制通过ascendc_fatbin_library生成include(${ASCENDC_CMAKE_DIR}/ascendc.cmake) ascendc_fatbin_library(custom_kernels kernel/custom.cpp)该指令把kernel/custom.cpp中的三个核函数打包为名为custom_kernels的独立算子二进制fatbin。Host 侧可执行程序binary_get_function_count链接acl_rt动态库并额外链接dl与pthread。构建要求 C17 标准编译选项包含-Wall -Wextra -Werror与-D_GLIBCXX_USE_CXX11_ABI0。编译运行1. 获取样例并进入目录将仓库代码下载至已安装 CANN 软件的环境后切换到样例目录cd ${git_clone_path}/example/2_advanced_features/kernel/3_binary_get_function_count2. 设置环境变量# ${install_root} 替换为 CANN 安装根目录默认安装在 /usr/local/Ascend 目录 source ${install_root}/cann/set_env.sh # 自动识别 SOC_VERSION 和 ASCENDC_CMAKE_DIR source ${git_clone_path}/example/set_sample_env.shset_sample_env.sh 会自动完成三项探测CANN 安装路径通过 resolve_cann_env.sh 在常见安装路径中查找并 sourceset_env.sh得到ASCEND_INSTALL_PATH/ASCEND_HOME_PATHSOC_VERSION样例内置的 get_soc_version 辅助程序 通过aclrtGetSocName查询真实芯片型号并输出 SOC 名称字符串ASCENDC_CMAKE_DIR按aarch64-linux/x86_64-linux等架构目录依次探测tikcpp/ascendc_kernel_cmake以ascendc.cmake是否存在作为判定依据。也可以手工指定关键变量例如export ASCEND_INSTALL_PATH/usr/local/Ascend/cann export SOC_VERSIONAscend910B1需要注意运行本样例除了 Runtime 组件外还必须具备 AscendC 编译工具链即ascendc.cmake。仅安装 Runtime 包时可能不存在该文件run.sh 在探测失败时会给出明确报错提示。3. 执行样例bash run.shrun.sh的执行逻辑值得关注同样适用于其他 CANN 样例若ASCEND_INSTALL_PATH/ASCEND_HOME_PATH、SOC_VERSION、ASCENDC_CMAKE_DIR任一为空自动 sourceset_sample_env.sh完成探测根据uname -m确定CANN_ARCH_DIRaarch64-linux或x86_64-linux使用 CMake 构建并安装到out目录算子二进制输出到out/fatbin/custom_kernels/custom_kernels.oHost 程序输出到out/bin/binary_get_function_count将runtime/lib64、lib64、架构目录下lib64依次加入LD_LIBRARY_PATH以算子二进制路径为参数运行 Host 程序。整个流程从加载二进制到校验结果完全自动化成功时以aclrtBinaryGetFunctionCount sample PASSED收尾。核心流程与接口调用链主程序 main.cpp 的流程可拆解为七个阶段下面结合关键代码片段逐段剖析。阶段一Runtime 初始化int32_t InitializeRuntime(RuntimeResources* runtime) { CHECK_ERROR(aclInit(nullptr)); runtime-aclInitialized true; CHECK_ERROR(aclrtSetDevice(runtime-deviceId)); runtime-deviceSet true; CHECK_ERROR(aclrtCreateStream(runtime-stream)); runtime-streamCreated true; return 0; }aclInit(nullptr)以默认配置完成 Runtime 初始化aclrtSetDevice(0)指定 Device 0 作为运算设备aclrtCreateStream创建任务下发所需的 Stream。值得注意的是RuntimeResources结构体为每个资源都保留了状态位aclInitialized、deviceSet、streamCreated、binLoaded用于在后续异常或收尾阶段精准释放已成功创建的资源避免重复释放或漏释放。阶段二加载算子二进制int32_t LoadBinary(const char* binaryPath, RuntimeResources* runtime) { CHECK_ERROR(aclrtBinaryLoadFromFile(binaryPath, nullptr, runtime-binHandle)); runtime-binLoaded true; return 0; }aclrtBinaryLoadFromFile的原型定义于 include/external/acl/acl_rt.hACL_FUNC_VISIBILITY aclError aclrtBinaryLoadFromFile( const char* binPath, aclrtBinaryLoadOptions* options, aclrtBinHandle* binHandle);binPath算子二进制文件路径本样例即custom_kernels.ooptions加载选项本样例传nullptrbinHandle输出参数返回二进制句柄供后续所有查询与卸载接口使用。该接口完成 Host 侧文件加载与解析真正把算子二进制关联到当前 Context 对应 Device 的动作会在首次通过该句柄获取函数aclrtBinaryGetFunction或查询函数数量时触发相关机制可对照 2_binary_enumerate_functions 样例 中关于aclrtBinaryEnumerateFunctions首次调用时拷贝到 Device 的说明。阶段三查询核函数总数本样例核心int32_t QueryFunctionCount(const RuntimeResources runtime, uint32_t functionCount) { CHECK_ACL(aclrtBinaryGetFunctionCount(runtime.binHandle, functionCount)); INFO_LOG(aclrtBinaryGetFunctionCount returned %u functions, functionCount); if (functionCount ! kExpectedFunctionCount) { ERROR_LOG(Expected %u functions, but got %u, kExpectedFunctionCount, functionCount); return -1; } return 0; }aclrtBinaryGetFunctionCount在头文件中的声明为见 include/external/acl/acl_rt.h/** * brief Get the total number of kernel functions in the binary * param [in] binHandle bin handle * param [out] count total number of kernel functions * retval ACL_SUCCESS The function is successfully executed. * retval OtherValues Failure */ ACL_FUNC_VISIBILITY aclError aclrtBinaryGetFunctionCount( const aclrtBinHandle binHandle, uint32_t *count);binHandle阶段二加载得到的二进制句柄count输出参数接收二进制内核函数总数。样例将返回值与预期值kExpectedFunctionCount 3比对一旦不一致立即报错退出。这种查询即校验的写法能让使用者在正式下发任务前就确认二进制内容与预期一致是实际工程中值得借鉴的防御性做法。阶段四按名称获取核函数句柄int32_t GetFunctionHandlesByName(const RuntimeResources runtime, ExpectedFunction* functions, uint32_t count) { for (uint32_t i 0U; i count; i) { CHECK_ACL(aclrtBinaryGetFunction(runtime.binHandle, functions[i].name, functions[i].handle)); } return 0; }aclrtBinaryGetFunction依据核函数名称从二进制中查找对应句柄见 include/external/acl/acl_rt.hACL_FUNC_VISIBILITY aclError aclrtBinaryGetFunction( const aclrtBinHandle binHandle, const char* kernelName, aclrtFuncHandle* funcHandle);样例预先定义了一张期望函数表将名称与期望输出值绑定为后续下发与校验提供数据驱动的基础ExpectedFunction expectedFunctions[kExpectedFunctionCount] { {add_custom, 3.0f, nullptr}, {sub_custom, -1.0f, nullptr}, {mul_custom, 2.0f, nullptr}, };输入为x1.0、y2.0FP16因此三个核函数的期望结果分别为3.0、-1.0和2.0。阶段五查询函数名称与地址int32_t QueryFunctionInfo(ExpectedFunction* functions, uint32_t count) { for (uint32_t i 0U; i count; i) { char functionName[kFunctionNameLength] {}; void* aicAddress nullptr; void* aivAddress nullptr; CHECK_ACL(aclrtGetFunctionName(functions[i].handle, kFunctionNameLength, functionName)); CHECK_ACL(aclrtGetFunctionAddr(functions[i].handle, aicAddress, aivAddress)); INFO_LOG( function[%u]: name%s, handle%p, aic%p, aiv%p, i, functionName, functions[i].handle, aicAddress, aivAddress); } return 0; }两个查询接口的原型分别位于 acl_rt.h 与 acl_rt.hACL_FUNC_VISIBILITY aclError aclrtGetFunctionAddr(aclrtFuncHandle funcHandle, void** aicAddr, void** aivAddr); ACL_FUNC_VISIBILITY aclError aclrtGetFunctionName(aclrtFuncHandle funcHandle, uint32_t maxLen, char* name);aclrtGetFunctionName以固定长度缓冲区kFunctionNameLength 128回填函数名用于校验按名查找结果是否正确aclrtGetFunctionAddr同时输出aicAI Core 向量/标量入口与aivAI Vector两个 Device 侧算子起始地址可用于底层调试或与驱动层地址对接。阶段六数据准备与任务下发数据准备采用 Host/Device 双层缓冲先用aclrtMallocHost分配三块 Host 内存x、y、z用aclrtMalloc(..., ACL_MEM_MALLOC_HUGE_FIRST)分配对应的 Device 内存再通过aclrtMemcpy(..., ACL_MEMCPY_HOST_TO_DEVICE)将输入拷入 DeviceCHECK_ERROR(aclrtMalloc(buffers-xDevice, dataSize, ACL_MEM_MALLOC_HUGE_FIRST)); ... CHECK_ERROR(aclrtMemcpy(buffers.xDevice, dataSize, buffers.xHost, dataSize, ACL_MEMCPY_HOST_TO_DEVICE));任务下发使用aclrtLaunchKernelWithHostArgsHost 侧参数在调用时由 Runtime 自动拷贝到 Devicevoid* args[] {buffers.xDevice, buffers.yDevice, buffers.zDevice}; CHECK_ACL(aclrtLaunchKernelWithHostArgs(handle, kBlockDim, stream, nullptr, args, sizeof(args), nullptr, 0U)); CHECK_ACL(aclrtSynchronizeStream(stream)); CHECK_ERROR(aclrtMemcpy(buffers.zHost, dataSize, buffers.zDevice, dataSize, ACL_MEMCPY_DEVICE_TO_HOST));该接口声明见 include/external/acl/acl_rt.hACL_FUNC_VISIBILITY aclError aclrtLaunchKernelWithHostArgs( aclrtFuncHandle funcHandle, uint32_t numBlocks, aclrtStream stream, aclrtLaunchKernelCfg* cfg, void* hostArgs, size_t argsSize, aclrtPlaceHolderInfo* placeHolderArray, size_t placeHolderNum);funcHandle阶段四获取的核函数句柄numBlocksblock 维度本样例为kBlockDim 8与核函数侧USE_CORE_NUM 8呼应hostArgs/argsSize以参数数组形式组织的入参三个 Device 指针cfg、placeHolderArray、placeHolderNum本样例均传nullptr/0。下发后调用aclrtSynchronizeStream阻塞等待任务完成再通过ACL_MEMCPY_DEVICE_TO_HOST将结果拷回 Host 校验。阶段七结果校验与资源释放校验逻辑对每一个输出元素做全量比对样例自行实现了Fp16ToFloat将 IEEE 754 half 手工转换为 float逐元素确认z[i]等于期望值任一元素不符即报错三个核函数全部通过后打印result verified。收尾阶段调用aclrtBinaryUnLoad卸载算子二进制、aclrtDestroyStreamForce强制销毁 Stream、aclrtResetDeviceForce强制复位 Device 并回收资源、aclFinalize完成去初始化对应的内存释放使用aclrtFree/aclrtFreeHost。整个流程演示了 CANN Runtime 中分配-使用-回收的完整资源生命周期管理。核函数侧实现解析三个核函数位于 kernel/custom.cpp采用 AscendC 编程模型编写。代码通过一个模板类按操作类型复用同一套搬入-计算-搬出流水template BinaryOperation operation class KernelBinary { public: __aicore__ inline void Init(GM_ADDR x, GM_ADDR y, GM_ADDR z) { ... } __aicore__ inline void Process() { ... } private: AscendC::TPipe pipe; AscendC::TQueAscendC::QuePosition::VECIN, BUFFER_NUM inQueueX; AscendC::TQueAscendC::QuePosition::VECIN, BUFFER_NUM inQueueY; AscendC::TQueAscendC::QuePosition::VECOUT, BUFFER_NUM outQueueZ; AscendC::GlobalTensorhalf xGm; AscendC::GlobalTensorhalf yGm; AscendC::GlobalTensorhalf zGm; };数据规模TOTAL_LENGTH 8 * 2048USE_CORE_NUM 8每个核处理BLOCK_LENGTH 2048个元素流水参数TILE_NUM 8、BUFFER_NUM 2即每个核心把数据切分为 16 个 Tile使用双缓冲隐藏搬运延迟Process内循环执行CopyInDataCopy全局到本地、ComputeAdd/Sub/Mul、CopyOutDataCopy本地回写全局。三个核函数以extern C __global__ __aicore__导出符号名称分别为add_custom、sub_custom和mul_custom这正是 Host 侧按名称查找的依据extern C __global__ __aicore__ void add_custom(GM_ADDR x, GM_ADDR y, GM_ADDR z) { ... } extern C __global__ __aicore__ void sub_custom(GM_ADDR x, GM_ADDR y, GM_ADDR z) { ... } extern C __global__ __aicore__ void mul_custom(GM_ADDR x, GM_ADDR y, GM_ADDR z) { ... }由于三者编译进同一个 fatbinaclrtBinaryGetFunctionCount才能一次性统计出总数为 3。这也解释了样例的查询-按名取句柄设计当二进制包含多个核函数时先查总数确认规模再按需逐个取句柄可避免对二进制内容的先验假设。CANN RUNTIME API 全景该样例覆盖了 CANN Runtime 五类核心接口汇总如下类别接口作用初始化aclInit完成 Runtime 初始化配置初始化aclFinalize去初始化释放全局资源Device 管理aclrtSetDevice指定用于运算的 DeviceDevice 管理aclrtResetDeviceForce强制复位当前运算的 Device回收 Device 上资源Stream 管理aclrtCreateStream创建 StreamStream 管理aclrtSynchronizeStream阻塞等待 Stream 上任务执行完成Stream 管理aclrtDestroyStreamForce强制销毁 Stream内存管理aclrtMallocHost分配 Host 内存内存管理aclrtMalloc分配 Device 内存内存管理aclrtMemcpy执行 Host 与 Device 之间的内存拷贝内存管理aclrtFreeHost释放 Host 内存内存管理aclrtFree释放 Device 内存二进制管理aclrtBinaryLoadFromFile从文件加载并解析算子二进制文件输出 binHandle二进制管理aclrtBinaryGetFunctionCount获取算子二进制中核函数句柄的总数二进制管理aclrtBinaryGetFunction根据核函数名称获取对应的核函数句柄二进制管理aclrtGetFunctionName根据核函数句柄获取核函数名称二进制管理aclrtGetFunctionAddr根据核函数句柄获取 Device 侧算子起始地址二进制管理aclrtLaunchKernelWithHostArgs下发 Kernel 计算任务Host 侧参数在调用时自动拷贝到 Device二进制管理aclrtBinaryUnLoad卸载算子二进制文件释放相关资源示例输出解读运行成功时输出类似如下内容...表示与环境相关的动态值[INFO] ASCEND_HOME_PATH... [INFO] ASCENDC_CMAKE_DIR... [INFO] SOC_VERSION... Configuring CMake... Building... [INFO] Querying the number of kernel functions in the binary [INFO] aclrtBinaryGetFunctionCount returned 3 functions [INFO] Getting function handles by name [INFO] Querying function names and addresses [INFO] function[0]: nameadd_custom, handle..., aic..., aiv... [INFO] function[1]: namesub_custom, handle..., aic..., aiv... [INFO] function[2]: namemul_custom, handle..., aic..., aiv... [INFO] Preparing input data [INFO] Launching kernels and verifying results [INFO] add_custom result verified [INFO] sub_custom result verified [INFO] mul_custom result verified [INFO] aclrtBinaryGetFunctionCount sample PASSED输出中的关键信息对应前文的每个阶段returned 3 functions验证了aclrtBinaryGetFunctionCount的查询结果三个function[i]行展示了按名获取句柄后回查的函数名、句柄与 aic/aiv 地址三个result verified表明各核函数均通过全量元素校验最后的PASSED表示整个样例运行成功。延伸阅读若希望无需预先知道函数名、直接遍历二进制内全部核函数可对照阅读姊妹样例 2_binary_enumerate_functions它演示了aclrtBinaryEnumerateFunctions的枚举式用法若希望进一步了解带资源限制、blocking 模式等更丰富的 Kernel 下发方式可参考仓库内 kernel 系列样例所有二进制管理接口的完整声明与参数说明均可查阅头文件 include/external/acl/acl_rt.h。【免费下载链接】runtime本项目提供CANN运行时组件和维测功能组件。项目地址: https://gitcode.com/cann/runtime创作声明：本文部分内容由AI辅助生成（AIGC），仅供参考