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 主要有两种下发方式：

需要特别注意的是：两种方式下，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_cast<size_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_custom<<<numBlocks, 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_custom<<<numBlocks, 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 句柄。
• Function：Binary 内部的具体可执行 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_cast<size_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 接口族选型

从源码实现看，这些接口在 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 训练/推理系列产品。

编译运行步骤

1. 切换到样例目录：

cd ${git_clone_path}/example/2_advanced_features/kernel/0_launch_kernel

2. 设置环境变量：

# ${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。

3. 运行样例（mode可选simple或placeholder，不指定时默认为simple）：

bash run.sh -r simple

simple模式中，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 接口调用顺序如下（也即非混合编程方式的标准操作序列）：

1. 初始化：aclInit→aclrtSetDevice→aclrtCreateStream；
2. 内存准备：aclrtMallocHost/aclrtMalloc申请 Host 与 Device 内存；
3. 数据搬运：aclrtMemcpy将输入从 Host 拷贝到 Device；
4. Kernel 加载与执行：aclrtBinaryLoadFromFile加载二进制 →aclrtBinaryGetFunction获取函数句柄 →aclrtKernelArgsInit初始化参数列表 →aclrtKernelArgsAppend追加参数（placeholder 模式追加aclrtKernelArgsAppendPlaceHolder+aclrtKernelArgsGetPlaceHolderBuffer）→aclrtKernelArgsFinalize→aclrtLaunchKernelWithConfig下发 →aclrtSynchronizeStream等待完成；
5. 结果回读：aclrtMemcpy将输出从 Device 拷贝到 Host；
6. 资源回收：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 correct

run.sh通过scripts/gen_data.py生成输入数据与 golden 参考结果，Kernel 执行输出写入output/output_z.bin，最后由scripts/verify_result.py对比误差比例（error ratio: 0.0000表示结果完全一致），并以[SUCCESS] result correct判定通过。

选型建议与注意事项

综合原文档与样例实践，可归纳出以下结论：

1. 能用 aclnn 就不用 LaunchKernel：若业务只是调用 CANN 已提供的算子，应优先使用 aclnn 算子接口，无需关心 Kernel 加载细节；
2. 混合编程适合 Kernel 与 Host 绑定发布：Kernel 集合在编译期已确定、代码追求简洁可读时，<<< >>>语法是最优解，且无需手动管理 Binary 生命周期；
3. 非混合编程适合独立发布与动态加载：需要 Kernel 二进制独立发布、按需选择（如按 tiling 选择不同实现）、延迟加载，或需要 placeholder、参数更新、任务属性配置等精细控制能力时，选择 aclrtBinary + LaunchKernel 接口族；
4. 参数组织方式决定了控制粒度：Device 参数区适合参数已在 Device 侧组装好的场景；Host 参数区与参数数组适合 Host 侧连续组织参数；aclrtArgsHandle + placeholder 则把参数布局交给 Runtime 管理，可避免为小尺寸 tiling 参数单独申请 Device 内存；
5. 别忘了同步与卸载：无论哪种方式，任务下发都是异步的，必须通过 Stream/Device/Event 同步等待；使用非混合编程时，还应在任务结束后调用aclrtBinaryUnLoad释放 Binary 资源，避免资源泄漏。

如需进一步理解底层实现，可从 src/acl/aclrt_impl/kernel.cpp（LaunchKernel/Binary/Function 接口实现）与 src/runtime/api/impl/api_impl_kernel_args.cc（参数列表组装实现）入手继续深入。

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