CANN Runtime 自定义 Kernel 加载与执行全指南混合编程与非混合编程实践【免费下载链接】runtime本项目提供CANN运行时组件和维测功能组件。项目地址: https://gitcode.com/cann/runtime本文基于 CANN/runtime 仓库的开发者指南与可运行样例系统讲解在昇腾平台上加载并执行自定义 Kernel 的两种方式混合编程 语法与非混合编程aclrtBinary / aclrtLaunchKernel 系列 Runtime 接口。读完本文你将掌握两种方式的差异与选型、Ascend C 工程的 CMake 构建配置、算子二进制的加载/卸载与函数句柄获取、Device/Host/placeholder 三种参数组织方式以及完整可运行的样例验证流程。何时需要关注 Kernel 加载与执行在 CANN 生态中普通用户通常不需要直接操作 Kernel。如果业务只是调用 CANN 已提供的算子能力应优先调用 aclnn 算子接口例如 aclnnAdd 等这些接口内部封装了算子选择、参数处理、Kernel 加载与执行全流程。只有在以下场景才需要直接使用 Kernel 加载与执行接口自己使用 Ascend C 编写了自定义算子 Kernel需要直接控制 Kernel 二进制加载、参数组装和任务下发过程需要将算子 Kernel 以独立二进制形式发布、动态选择或延迟加载。两种下发方式总览自定义 Kernel 主要有两种下发方式对比项混合编程 非混合编程LaunchKernel 接口Kernel 引用方式Host 代码直接引用 Kernel 函数符号例如add_custom...(...)Host 代码通过 Kernel 名称字符串获取 aclrtFuncHandle例如aclrtBinaryGetFunction(bin, add_custom, func)编译方式Kernel 源码参与 Host 工程构建通常通过 Ascend C CMake 能力生成可链接的 Kernel 库并与 Host 可执行文件一起链接Kernel 源码单独编译为算子二进制文件如*.o或 fatbinHost 程序单独编译在运行时通过 Runtime 接口加载该二进制生成代码构建系统生成 Host 侧可调用的 Kernel 启动桩、注册信息和设备侧代码绑定关系使 语法能够直接下发任务不生成可直接调用的 Kernel 启动桩Host 侧只保存二进制路径、Kernel 名称和 Runtime 句柄需要手动加载 Binary、获取 Function 并组装参数运行时加载Kernel 二进制通常随 Host 程序或链接库注册首次下发时由生成代码完成注册和加载用户显式调用 aclrtBinaryLoadFromFile 或 aclrtBinaryLoadFromData 加载 Binary结束时调用 aclrtBinaryUnLoad 卸载参数组织以函数调用形式传参代码简洁、可读性好可使用 Device 参数区、Host 参数区、aclrtArgsHandle 参数列表、placeholder 等方式组织参数控制粒度更细适用场景Kernel 与 Host 程序一起开发、一起发布Kernel 集合在编译期已确定Kernel 二进制需要独立发布、动态选择、延迟加载或需要使用 Runtime 参数组装、placeholder、任务属性配置等能力需要特别注意的是两种方式下Kernel 任务下发后相对 Host 线程都是异步执行。Host 线程如需等待 Kernel 执行完成必须调用 aclrtSynchronizeStream、aclrtSynchronizeDevice 或使用 Event 等同步机制。混合编程方式 直接下发混合编程方式中Host 代码可以直接调用 Kernel 函数并使用 语法指定任务的 block 数量、动态共享内存参数和 Stream。该方式代码最简洁适合 Kernel 与 Host 程序绑定发布的场景。构建配置编译时Kernel 源码作为工程的一部分参与构建。使用 Ascend C 提供的 CMake 能力可将 Kernel 源码编译为静态库再链接到 Host 可执行文件include(${ASCENDC_CMAKE_DIR}/ascendc.cmake) ascendc_library(kernels STATIC kernel_print.cpp) add_executable(main main.cpp) target_link_libraries(main PRIVATE kernels ${ASCEND_CANN_PACKAGE_PATH}/lib64/libacl_rt.so)其中ASCENDC_CMAKE_DIR指向 Ascend C CMake 模块所在目录ascendc_library负责将 Kernel 源码编译为可链接的 Kernel 库ASCEND_CANN_PACKAGE_PATH指向 CANN 安装路径。上述方式生成的 Host 侧代码可以直接调用 Kernel 启动函数不需要显式调用 aclrtBinaryLoadFromFile、aclrtBinaryGetFunction 等接口。Host 侧关键代码以下示例不可直接拷贝编译运行仅用于理解流程对应 Kernel 设备侧代码见 example/kernel_func/add_custom.cpp// Device code extern C __global__ __aicore__ void add_custom(GM_ADDR x, GM_ADDR y, GM_ADDR z) { KernelAdd op; op.Init(x, y, z); op.Process(); } int main() { int64_t n ...; size_t size static_castsize_t(n) * sizeof(uint64_t); aclInit(nullptr); aclrtSetDevice(0); aclrtStream stream nullptr; aclrtCreateStream(stream); void *hX nullptr; void *hY nullptr; void *hZ nullptr; aclrtMallocHost(hX, size); aclrtMallocHost(hY, size); aclrtMallocHost(hZ, size); // 初始化输入数据。 ... void *dX nullptr; void *dY nullptr; void *dZ nullptr; aclrtMalloc(dX, size, ACL_MEM_MALLOC_HUGE_FIRST); aclrtMalloc(dY, size, ACL_MEM_MALLOC_HUGE_FIRST); aclrtMalloc(dZ, size, ACL_MEM_MALLOC_HUGE_FIRST); aclrtMemcpy(dX, size, hX, size, ACL_MEMCPY_HOST_TO_DEVICE); aclrtMemcpy(dY, size, hY, size, ACL_MEMCPY_HOST_TO_DEVICE); uint32_t numBlocks 48; add_customnumBlocks, nullptr, stream(dX, dY, dZ); aclrtSynchronizeStream(stream); aclrtMemcpy(hZ, size, dZ, size, ACL_MEMCPY_DEVICE_TO_HOST); ... }从代码可以看到混合编程方式的核心特点aclInit/aclrtSetDevice完成 ACL 初始化与 Device 选择aclrtMallocHost与aclrtMalloc分别申请 Host 与 Device 内存aclrtMemcpy完成 H2D/D2H 数据传输add_customnumBlocks, nullptr, stream(dX, dY, dZ)直接以函数调用形式下发nullptr表示不指定动态共享内存aclrtSynchronizeStream阻塞等待 Kernel 执行完成。非混合编程方式LaunchKernel 接口族非混合编程方式中Kernel 设备侧代码先编译成独立算子二进制Host 程序运行时显式加载该二进制获取 Kernel 函数句柄组装参数后调用 LaunchKernel 接口下发任务。构建配置Kernel 源码与 Host 程序可以分开构建。使用ascendc_fatbin_library生成算子二进制文件再由 Host 程序运行时加载。样例 example/2_advanced_features/kernel/0_launch_kernel/CMakeLists.txt 给出了完整写法include(${ASCENDC_CMAKE_DIR}/ascendc.cmake) ascendc_fatbin_library(ascendc_kernels_simple add_custom.cpp) add_executable(ascendc_kernels_bbit main.cpp) target_link_libraries(ascendc_kernels_bbit PRIVATE ${ASCEND_CANN_PACKAGE_PATH}/lib64/libacl_rt.so)生成的算子二进制文件在运行时通过路径加载例如./out/fatbin/ascendc_kernels_simple/ascendc_kernels_simple.o。这种模式下 Host 代码不直接引用 Kernel 函数符号而是通过 Binary 和 Function 句柄操作 Kernel。三个核心概念Binary动态加载的代码容器单元包含编译后的 Kernel 代码、全局变量等。通过aclrtBinaryLoadFromFile或aclrtBinaryLoadFromData加载算子二进制并获得 Binary 句柄。FunctionBinary 内部的具体可执行 Kernel 入口。通过aclrtBinaryGetFunction或aclrtBinaryGetFunctionByEntry获取 Function 句柄。参数列表LaunchKernel 接口需要获取 Kernel 参数。参数可放在 Device 内存、Host 内存或 aclrtArgsHandle 参数列表中也可以使用 placeholder 让 Runtime 在 Launch 时完成小块参数数据的搬运。Host 侧关键代码以下示例不可直接拷贝编译运行仅用于理解流程// Device code编译为独立算子二进制文件。 extern C __global__ __aicore__ void add_custom(GM_ADDR x, GM_ADDR y, GM_ADDR z) { KernelAdd op; op.Init(x, y, z); op.Process(); } int main() { int64_t n ...; size_t size static_castsize_t(n) * sizeof(uint64_t); aclInit(nullptr); aclrtSetDevice(0); aclrtStream stream nullptr; aclrtCreateStream(stream); // 加载算子二进制并获取Kernel函数句柄。 aclrtBinHandle bin nullptr; aclrtBinaryLoadFromFile(add_custom.o, nullptr, bin); aclrtFuncHandle func nullptr; aclrtBinaryGetFunction(bin, add_custom, func); void *hX nullptr; void *hY nullptr; void *hZ nullptr; aclrtMallocHost(hX, size); aclrtMallocHost(hY, size); aclrtMallocHost(hZ, size); // 初始化输入数据。 ... void *dX nullptr; void *dY nullptr; void *dZ nullptr; aclrtMalloc(dX, size, ACL_MEM_MALLOC_HUGE_FIRST); aclrtMalloc(dY, size, ACL_MEM_MALLOC_HUGE_FIRST); aclrtMalloc(dZ, size, ACL_MEM_MALLOC_HUGE_FIRST); aclrtMemcpy(dX, size, hX, size, ACL_MEMCPY_HOST_TO_DEVICE); aclrtMemcpy(dY, size, hY, size, ACL_MEMCPY_HOST_TO_DEVICE); uint32_t numBlocks 48; void *args[] {dX, dY, dZ}; size_t argsSize sizeof(args); aclrtLaunchKernelWithHostArgs(func, numBlocks, stream, nullptr, args, argsSize, nullptr, 0); aclrtSynchronizeStream(stream); aclrtMemcpy(hZ, size, dZ, size, ACL_MEMCPY_DEVICE_TO_HOST); aclrtBinaryUnLoad(bin); ... }这段代码与混合编程方式的差别在于不再直接引用 Kernel 符号而是通过aclrtBinaryLoadFromFile加载二进制、aclrtBinaryGetFunction按名称获取 Function 句柄参数以void* args[]数组组织后由aclrtLaunchKernelWithHostArgs下发最后调用aclrtBinaryUnLoad卸载。LaunchKernel 接口族选型接口参数来源适用场景aclrtLaunchKernelDevice 内存中的完整参数区参数已经在 Device 侧组装完成不需要 Launch 配置aclrtLaunchKernelV2Device 内存中的完整参数区需要额外指定 Launch 配置aclrtLaunchKernelWithConfigaclrtArgsHandle 参数列表希望由 Runtime 管理参数布局或需要使用 placeholder、参数更新等能力aclrtLaunchKernelWithHostArgsHost 内存中的完整参数区参数在 Host 侧连续组织Launch 时由 Runtime 处理aclrtLaunchKernelWithArgsArrayHost 侧参数数组每个数组元素指向一个参数数据便于按参数数组形式组织调用从源码实现看这些接口在 src/acl/aclrt_impl/kernel.cpp 中均有对应实现aclrtBinaryLoadFromFileImpl、aclrtBinaryGetFunctionImpl、aclrtLaunchKernelWithConfigImpl、aclrtLaunchKernelV2Impl、aclrtLaunchKernelWithHostArgsImpl、aclrtLaunchKernelWithArgsArrayImpl等参数列表相关接口aclrtKernelArgsInitImpl、aclrtKernelArgsAppendImpl、aclrtKernelArgsAppendPlaceHolderImpl等则集中在 src/runtime/api/impl/api_impl_kernel_args.cc可以作为阅读底层实现时的入口。实战样例0_launch_kernel 完整流程仓库中的 example/2_advanced_features/kernel/0_launch_kernel 是上述接口的完整可运行样例覆盖二进制加载、核函数句柄获取、参数组装、任务下发、Stream 同步和结果校验支持simple与placeholder两种参数组织模式。该样例支持 Ascend 950PR/Ascend 950DT、Atlas A3 与 Atlas A2 训练/推理系列产品。编译运行步骤切换到样例目录cd ${git_clone_path}/example/2_advanced_features/kernel/0_launch_kernel设置环境变量# ${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.sh样例的数据生成与结果校验依赖numpy执行run.sh前请确保 Python 环境已安装numpy。运行样例mode可选simple或placeholder不指定时默认为simplebash run.sh -r simplesimple模式中Kernel 指针类型参数使用用户提前申请并拷贝数据后的 Device 内存地址placeholder模式中placeholder 参数对应的数据由 Runtime 在 Kernel Launch 时传输到 Device 侧。两种模式的参数组织差异两种模式共用同一套 aclrtArgsHandle 组装框架见 main.cpp差异体现在是否追加 placeholder 参数simple 模式三个指针参数x、y、z 的 Device 地址通过aclrtKernelArgsAppend逐项追加到参数列表参数值就是用户提前aclrtMalloc并aclrtMemcpy数据后的 Device 地址placeholder 模式前三个指针参数同上另外通过aclrtKernelArgsAppendPlaceHolder追加两个占位参数再调用aclrtKernelArgsGetPlaceHolderBuffer获取 Runtime 分配的 Host 缓冲区并写入 tiling 数值TOTAL_LENGTH 8 * 2048、TILE_NUM 8。Kernel Launch 时 Runtime 会自动将这些小块参数数据搬运到 Device 侧用户无需自行申请 Device 内存和拷贝。两种模式对应的 Kernel 源码也体现了这种差异simple 模式 Kernel add_custom.cpp 只有三个GM_ADDR指针参数数据规模由编译期常量TOTAL_LENGTH、TILE_NUM决定placeholder 模式 Kernel add_custom_tiling.cpp 除三个指针参数外还接收__gm__ int32_t* tilingLength与__gm__ int32_t* tilingNum两个 tiling 参数由 Kernel 内部读取后再计算tileLength等分块信息——tiling 信息从编译期常量变成了运行时参数。两种模式的参数组装顺序都是aclrtKernelArgsInit初始化参数列表 → 追加指针参数必要时追加 placeholder 并填充 Host 缓冲区→aclrtKernelArgsFinalize标识参数组装完毕 →aclrtLaunchKernelWithConfig下发任务。完整的接口调用链样例涉及的关键 Runtime 接口调用顺序如下也即非混合编程方式的标准操作序列初始化aclInit→aclrtSetDevice→aclrtCreateStream内存准备aclrtMallocHost/aclrtMalloc申请 Host 与 Device 内存数据搬运aclrtMemcpy将输入从 Host 拷贝到 DeviceKernel 加载与执行aclrtBinaryLoadFromFile加载二进制 →aclrtBinaryGetFunction获取函数句柄 →aclrtKernelArgsInit初始化参数列表 →aclrtKernelArgsAppend追加参数placeholder 模式追加aclrtKernelArgsAppendPlaceHolderaclrtKernelArgsGetPlaceHolderBuffer→aclrtKernelArgsFinalize→aclrtLaunchKernelWithConfig下发 →aclrtSynchronizeStream等待完成结果回读aclrtMemcpy将输出从 Device 拷贝到 Host资源回收aclrtBinaryUnLoad卸载二进制 →aclrtFreeHost/aclrtFree释放内存 →aclrtDestroyStreamForce销毁 Stream →aclrtResetDeviceForce复位 Device →aclFinalize去初始化。运行结果验证样例输出形如Configuring CMake... Building... ... [INFO] Kernel launch sample runs in simple mode. [INFO] Run the launch_kernel sample successfully. ... output/output_z.bin ... output/golden.bin error ratio: 0.0000, tolerance: 0.0010 [SUCCESS] result correctrun.sh通过scripts/gen_data.py生成输入数据与 golden 参考结果Kernel 执行输出写入output/output_z.bin最后由scripts/verify_result.py对比误差比例error ratio: 0.0000表示结果完全一致并以[SUCCESS] result correct判定通过。选型建议与注意事项综合原文档与样例实践可归纳出以下结论能用 aclnn 就不用 LaunchKernel若业务只是调用 CANN 已提供的算子应优先使用 aclnn 算子接口无需关心 Kernel 加载细节混合编程适合 Kernel 与 Host 绑定发布Kernel 集合在编译期已确定、代码追求简洁可读时 语法是最优解且无需手动管理 Binary 生命周期非混合编程适合独立发布与动态加载需要 Kernel 二进制独立发布、按需选择如按 tiling 选择不同实现、延迟加载或需要 placeholder、参数更新、任务属性配置等精细控制能力时选择 aclrtBinary LaunchKernel 接口族参数组织方式决定了控制粒度Device 参数区适合参数已在 Device 侧组装好的场景Host 参数区与参数数组适合 Host 侧连续组织参数aclrtArgsHandle placeholder 则把参数布局交给 Runtime 管理可避免为小尺寸 tiling 参数单独申请 Device 内存别忘了同步与卸载无论哪种方式任务下发都是异步的必须通过 Stream/Device/Event 同步等待使用非混合编程时还应在任务结束后调用aclrtBinaryUnLoad释放 Binary 资源避免资源泄漏。如需进一步理解底层实现可从 src/acl/aclrt_impl/kernel.cppLaunchKernel/Binary/Function 接口实现与 src/runtime/api/impl/api_impl_kernel_args.cc参数列表组装实现入手继续深入。【免费下载链接】runtime本项目提供CANN运行时组件和维测功能组件。项目地址: https://gitcode.com/cann/runtime创作声明：本文部分内容由AI辅助生成（AIGC），仅供参考 CANN Runtime 自定义 Kernel 加载与执行全指南：混合编程与非混合编程实践
CANN Runtime 自定义 Kernel 加载与执行全指南混合编程与非混合编程实践【免费下载链接】runtime本项目提供CANN运行时组件和维测功能组件。项目地址: https://gitcode.com/cann/runtime本文基于 CANN/runtime 仓库的开发者指南与可运行样例系统讲解在昇腾平台上加载并执行自定义 Kernel 的两种方式混合编程 语法与非混合编程aclrtBinary / aclrtLaunchKernel 系列 Runtime 接口。读完本文你将掌握两种方式的差异与选型、Ascend C 工程的 CMake 构建配置、算子二进制的加载/卸载与函数句柄获取、Device/Host/placeholder 三种参数组织方式以及完整可运行的样例验证流程。何时需要关注 Kernel 加载与执行在 CANN 生态中普通用户通常不需要直接操作 Kernel。如果业务只是调用 CANN 已提供的算子能力应优先调用 aclnn 算子接口例如 aclnnAdd 等这些接口内部封装了算子选择、参数处理、Kernel 加载与执行全流程。只有在以下场景才需要直接使用 Kernel 加载与执行接口自己使用 Ascend C 编写了自定义算子 Kernel需要直接控制 Kernel 二进制加载、参数组装和任务下发过程需要将算子 Kernel 以独立二进制形式发布、动态选择或延迟加载。两种下发方式总览自定义 Kernel 主要有两种下发方式对比项混合编程 非混合编程LaunchKernel 接口Kernel 引用方式Host 代码直接引用 Kernel 函数符号例如add_custom...(...)Host 代码通过 Kernel 名称字符串获取 aclrtFuncHandle例如aclrtBinaryGetFunction(bin, add_custom, func)编译方式Kernel 源码参与 Host 工程构建通常通过 Ascend C CMake 能力生成可链接的 Kernel 库并与 Host 可执行文件一起链接Kernel 源码单独编译为算子二进制文件如*.o或 fatbinHost 程序单独编译在运行时通过 Runtime 接口加载该二进制生成代码构建系统生成 Host 侧可调用的 Kernel 启动桩、注册信息和设备侧代码绑定关系使 语法能够直接下发任务不生成可直接调用的 Kernel 启动桩Host 侧只保存二进制路径、Kernel 名称和 Runtime 句柄需要手动加载 Binary、获取 Function 并组装参数运行时加载Kernel 二进制通常随 Host 程序或链接库注册首次下发时由生成代码完成注册和加载用户显式调用 aclrtBinaryLoadFromFile 或 aclrtBinaryLoadFromData 加载 Binary结束时调用 aclrtBinaryUnLoad 卸载参数组织以函数调用形式传参代码简洁、可读性好可使用 Device 参数区、Host 参数区、aclrtArgsHandle 参数列表、placeholder 等方式组织参数控制粒度更细适用场景Kernel 与 Host 程序一起开发、一起发布Kernel 集合在编译期已确定Kernel 二进制需要独立发布、动态选择、延迟加载或需要使用 Runtime 参数组装、placeholder、任务属性配置等能力需要特别注意的是两种方式下Kernel 任务下发后相对 Host 线程都是异步执行。Host 线程如需等待 Kernel 执行完成必须调用 aclrtSynchronizeStream、aclrtSynchronizeDevice 或使用 Event 等同步机制。混合编程方式 直接下发混合编程方式中Host 代码可以直接调用 Kernel 函数并使用 语法指定任务的 block 数量、动态共享内存参数和 Stream。该方式代码最简洁适合 Kernel 与 Host 程序绑定发布的场景。构建配置编译时Kernel 源码作为工程的一部分参与构建。使用 Ascend C 提供的 CMake 能力可将 Kernel 源码编译为静态库再链接到 Host 可执行文件include(${ASCENDC_CMAKE_DIR}/ascendc.cmake) ascendc_library(kernels STATIC kernel_print.cpp) add_executable(main main.cpp) target_link_libraries(main PRIVATE kernels ${ASCEND_CANN_PACKAGE_PATH}/lib64/libacl_rt.so)其中ASCENDC_CMAKE_DIR指向 Ascend C CMake 模块所在目录ascendc_library负责将 Kernel 源码编译为可链接的 Kernel 库ASCEND_CANN_PACKAGE_PATH指向 CANN 安装路径。上述方式生成的 Host 侧代码可以直接调用 Kernel 启动函数不需要显式调用 aclrtBinaryLoadFromFile、aclrtBinaryGetFunction 等接口。Host 侧关键代码以下示例不可直接拷贝编译运行仅用于理解流程对应 Kernel 设备侧代码见 example/kernel_func/add_custom.cpp// Device code extern C __global__ __aicore__ void add_custom(GM_ADDR x, GM_ADDR y, GM_ADDR z) { KernelAdd op; op.Init(x, y, z); op.Process(); } int main() { int64_t n ...; size_t size static_castsize_t(n) * sizeof(uint64_t); aclInit(nullptr); aclrtSetDevice(0); aclrtStream stream nullptr; aclrtCreateStream(stream); void *hX nullptr; void *hY nullptr; void *hZ nullptr; aclrtMallocHost(hX, size); aclrtMallocHost(hY, size); aclrtMallocHost(hZ, size); // 初始化输入数据。 ... void *dX nullptr; void *dY nullptr; void *dZ nullptr; aclrtMalloc(dX, size, ACL_MEM_MALLOC_HUGE_FIRST); aclrtMalloc(dY, size, ACL_MEM_MALLOC_HUGE_FIRST); aclrtMalloc(dZ, size, ACL_MEM_MALLOC_HUGE_FIRST); aclrtMemcpy(dX, size, hX, size, ACL_MEMCPY_HOST_TO_DEVICE); aclrtMemcpy(dY, size, hY, size, ACL_MEMCPY_HOST_TO_DEVICE); uint32_t numBlocks 48; add_customnumBlocks, nullptr, stream(dX, dY, dZ); aclrtSynchronizeStream(stream); aclrtMemcpy(hZ, size, dZ, size, ACL_MEMCPY_DEVICE_TO_HOST); ... }从代码可以看到混合编程方式的核心特点aclInit/aclrtSetDevice完成 ACL 初始化与 Device 选择aclrtMallocHost与aclrtMalloc分别申请 Host 与 Device 内存aclrtMemcpy完成 H2D/D2H 数据传输add_customnumBlocks, nullptr, stream(dX, dY, dZ)直接以函数调用形式下发nullptr表示不指定动态共享内存aclrtSynchronizeStream阻塞等待 Kernel 执行完成。非混合编程方式LaunchKernel 接口族非混合编程方式中Kernel 设备侧代码先编译成独立算子二进制Host 程序运行时显式加载该二进制获取 Kernel 函数句柄组装参数后调用 LaunchKernel 接口下发任务。构建配置Kernel 源码与 Host 程序可以分开构建。使用ascendc_fatbin_library生成算子二进制文件再由 Host 程序运行时加载。样例 example/2_advanced_features/kernel/0_launch_kernel/CMakeLists.txt 给出了完整写法include(${ASCENDC_CMAKE_DIR}/ascendc.cmake) ascendc_fatbin_library(ascendc_kernels_simple add_custom.cpp) add_executable(ascendc_kernels_bbit main.cpp) target_link_libraries(ascendc_kernels_bbit PRIVATE ${ASCEND_CANN_PACKAGE_PATH}/lib64/libacl_rt.so)生成的算子二进制文件在运行时通过路径加载例如./out/fatbin/ascendc_kernels_simple/ascendc_kernels_simple.o。这种模式下 Host 代码不直接引用 Kernel 函数符号而是通过 Binary 和 Function 句柄操作 Kernel。三个核心概念Binary动态加载的代码容器单元包含编译后的 Kernel 代码、全局变量等。通过aclrtBinaryLoadFromFile或aclrtBinaryLoadFromData加载算子二进制并获得 Binary 句柄。FunctionBinary 内部的具体可执行 Kernel 入口。通过aclrtBinaryGetFunction或aclrtBinaryGetFunctionByEntry获取 Function 句柄。参数列表LaunchKernel 接口需要获取 Kernel 参数。参数可放在 Device 内存、Host 内存或 aclrtArgsHandle 参数列表中也可以使用 placeholder 让 Runtime 在 Launch 时完成小块参数数据的搬运。Host 侧关键代码以下示例不可直接拷贝编译运行仅用于理解流程// Device code编译为独立算子二进制文件。 extern C __global__ __aicore__ void add_custom(GM_ADDR x, GM_ADDR y, GM_ADDR z) { KernelAdd op; op.Init(x, y, z); op.Process(); } int main() { int64_t n ...; size_t size static_castsize_t(n) * sizeof(uint64_t); aclInit(nullptr); aclrtSetDevice(0); aclrtStream stream nullptr; aclrtCreateStream(stream); // 加载算子二进制并获取Kernel函数句柄。 aclrtBinHandle bin nullptr; aclrtBinaryLoadFromFile(add_custom.o, nullptr, bin); aclrtFuncHandle func nullptr; aclrtBinaryGetFunction(bin, add_custom, func); void *hX nullptr; void *hY nullptr; void *hZ nullptr; aclrtMallocHost(hX, size); aclrtMallocHost(hY, size); aclrtMallocHost(hZ, size); // 初始化输入数据。 ... void *dX nullptr; void *dY nullptr; void *dZ nullptr; aclrtMalloc(dX, size, ACL_MEM_MALLOC_HUGE_FIRST); aclrtMalloc(dY, size, ACL_MEM_MALLOC_HUGE_FIRST); aclrtMalloc(dZ, size, ACL_MEM_MALLOC_HUGE_FIRST); aclrtMemcpy(dX, size, hX, size, ACL_MEMCPY_HOST_TO_DEVICE); aclrtMemcpy(dY, size, hY, size, ACL_MEMCPY_HOST_TO_DEVICE); uint32_t numBlocks 48; void *args[] {dX, dY, dZ}; size_t argsSize sizeof(args); aclrtLaunchKernelWithHostArgs(func, numBlocks, stream, nullptr, args, argsSize, nullptr, 0); aclrtSynchronizeStream(stream); aclrtMemcpy(hZ, size, dZ, size, ACL_MEMCPY_DEVICE_TO_HOST); aclrtBinaryUnLoad(bin); ... }这段代码与混合编程方式的差别在于不再直接引用 Kernel 符号而是通过aclrtBinaryLoadFromFile加载二进制、aclrtBinaryGetFunction按名称获取 Function 句柄参数以void* args[]数组组织后由aclrtLaunchKernelWithHostArgs下发最后调用aclrtBinaryUnLoad卸载。LaunchKernel 接口族选型接口参数来源适用场景aclrtLaunchKernelDevice 内存中的完整参数区参数已经在 Device 侧组装完成不需要 Launch 配置aclrtLaunchKernelV2Device 内存中的完整参数区需要额外指定 Launch 配置aclrtLaunchKernelWithConfigaclrtArgsHandle 参数列表希望由 Runtime 管理参数布局或需要使用 placeholder、参数更新等能力aclrtLaunchKernelWithHostArgsHost 内存中的完整参数区参数在 Host 侧连续组织Launch 时由 Runtime 处理aclrtLaunchKernelWithArgsArrayHost 侧参数数组每个数组元素指向一个参数数据便于按参数数组形式组织调用从源码实现看这些接口在 src/acl/aclrt_impl/kernel.cpp 中均有对应实现aclrtBinaryLoadFromFileImpl、aclrtBinaryGetFunctionImpl、aclrtLaunchKernelWithConfigImpl、aclrtLaunchKernelV2Impl、aclrtLaunchKernelWithHostArgsImpl、aclrtLaunchKernelWithArgsArrayImpl等参数列表相关接口aclrtKernelArgsInitImpl、aclrtKernelArgsAppendImpl、aclrtKernelArgsAppendPlaceHolderImpl等则集中在 src/runtime/api/impl/api_impl_kernel_args.cc可以作为阅读底层实现时的入口。实战样例0_launch_kernel 完整流程仓库中的 example/2_advanced_features/kernel/0_launch_kernel 是上述接口的完整可运行样例覆盖二进制加载、核函数句柄获取、参数组装、任务下发、Stream 同步和结果校验支持simple与placeholder两种参数组织模式。该样例支持 Ascend 950PR/Ascend 950DT、Atlas A3 与 Atlas A2 训练/推理系列产品。编译运行步骤切换到样例目录cd ${git_clone_path}/example/2_advanced_features/kernel/0_launch_kernel设置环境变量# ${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.sh样例的数据生成与结果校验依赖numpy执行run.sh前请确保 Python 环境已安装numpy。运行样例mode可选simple或placeholder不指定时默认为simplebash run.sh -r simplesimple模式中Kernel 指针类型参数使用用户提前申请并拷贝数据后的 Device 内存地址placeholder模式中placeholder 参数对应的数据由 Runtime 在 Kernel Launch 时传输到 Device 侧。两种模式的参数组织差异两种模式共用同一套 aclrtArgsHandle 组装框架见 main.cpp差异体现在是否追加 placeholder 参数simple 模式三个指针参数x、y、z 的 Device 地址通过aclrtKernelArgsAppend逐项追加到参数列表参数值就是用户提前aclrtMalloc并aclrtMemcpy数据后的 Device 地址placeholder 模式前三个指针参数同上另外通过aclrtKernelArgsAppendPlaceHolder追加两个占位参数再调用aclrtKernelArgsGetPlaceHolderBuffer获取 Runtime 分配的 Host 缓冲区并写入 tiling 数值TOTAL_LENGTH 8 * 2048、TILE_NUM 8。Kernel Launch 时 Runtime 会自动将这些小块参数数据搬运到 Device 侧用户无需自行申请 Device 内存和拷贝。两种模式对应的 Kernel 源码也体现了这种差异simple 模式 Kernel add_custom.cpp 只有三个GM_ADDR指针参数数据规模由编译期常量TOTAL_LENGTH、TILE_NUM决定placeholder 模式 Kernel add_custom_tiling.cpp 除三个指针参数外还接收__gm__ int32_t* tilingLength与__gm__ int32_t* tilingNum两个 tiling 参数由 Kernel 内部读取后再计算tileLength等分块信息——tiling 信息从编译期常量变成了运行时参数。两种模式的参数组装顺序都是aclrtKernelArgsInit初始化参数列表 → 追加指针参数必要时追加 placeholder 并填充 Host 缓冲区→aclrtKernelArgsFinalize标识参数组装完毕 →aclrtLaunchKernelWithConfig下发任务。完整的接口调用链样例涉及的关键 Runtime 接口调用顺序如下也即非混合编程方式的标准操作序列初始化aclInit→aclrtSetDevice→aclrtCreateStream内存准备aclrtMallocHost/aclrtMalloc申请 Host 与 Device 内存数据搬运aclrtMemcpy将输入从 Host 拷贝到 DeviceKernel 加载与执行aclrtBinaryLoadFromFile加载二进制 →aclrtBinaryGetFunction获取函数句柄 →aclrtKernelArgsInit初始化参数列表 →aclrtKernelArgsAppend追加参数placeholder 模式追加aclrtKernelArgsAppendPlaceHolderaclrtKernelArgsGetPlaceHolderBuffer→aclrtKernelArgsFinalize→aclrtLaunchKernelWithConfig下发 →aclrtSynchronizeStream等待完成结果回读aclrtMemcpy将输出从 Device 拷贝到 Host资源回收aclrtBinaryUnLoad卸载二进制 →aclrtFreeHost/aclrtFree释放内存 →aclrtDestroyStreamForce销毁 Stream →aclrtResetDeviceForce复位 Device →aclFinalize去初始化。运行结果验证样例输出形如Configuring CMake... Building... ... [INFO] Kernel launch sample runs in simple mode. [INFO] Run the launch_kernel sample successfully. ... output/output_z.bin ... output/golden.bin error ratio: 0.0000, tolerance: 0.0010 [SUCCESS] result correctrun.sh通过scripts/gen_data.py生成输入数据与 golden 参考结果Kernel 执行输出写入output/output_z.bin最后由scripts/verify_result.py对比误差比例error ratio: 0.0000表示结果完全一致并以[SUCCESS] result correct判定通过。选型建议与注意事项综合原文档与样例实践可归纳出以下结论能用 aclnn 就不用 LaunchKernel若业务只是调用 CANN 已提供的算子应优先使用 aclnn 算子接口无需关心 Kernel 加载细节混合编程适合 Kernel 与 Host 绑定发布Kernel 集合在编译期已确定、代码追求简洁可读时 语法是最优解且无需手动管理 Binary 生命周期非混合编程适合独立发布与动态加载需要 Kernel 二进制独立发布、按需选择如按 tiling 选择不同实现、延迟加载或需要 placeholder、参数更新、任务属性配置等精细控制能力时选择 aclrtBinary LaunchKernel 接口族参数组织方式决定了控制粒度Device 参数区适合参数已在 Device 侧组装好的场景Host 参数区与参数数组适合 Host 侧连续组织参数aclrtArgsHandle placeholder 则把参数布局交给 Runtime 管理可避免为小尺寸 tiling 参数单独申请 Device 内存别忘了同步与卸载无论哪种方式任务下发都是异步的必须通过 Stream/Device/Event 同步等待使用非混合编程时还应在任务结束后调用aclrtBinaryUnLoad释放 Binary 资源避免资源泄漏。如需进一步理解底层实现可从 src/acl/aclrt_impl/kernel.cppLaunchKernel/Binary/Function 接口实现与 src/runtime/api/impl/api_impl_kernel_args.cc参数列表组装实现入手继续深入。【免费下载链接】runtime本项目提供CANN运行时组件和维测功能组件。项目地址: https://gitcode.com/cann/runtime创作声明：本文部分内容由AI辅助生成（AIGC），仅供参考